[libclc] Move clcmacro.h to CLC library. NFC (#114845)

llvm · Nov 4, 2024 · d2d1b58 · d2d1b58
1 parent 8d023b7
commit d2d1b58
Show file tree

Hide file tree

Showing 95 changed files with 220 additions and 220 deletions.
diff --git a/libclc/amdgcn/lib/integer/popcount.cl b/libclc/amdgcn/lib/integer/popcount.cl
@@ -1,5 +1,5 @@
 #include <clc/clc.h>
-#include <utils.h>
+#include <clc/utils.h>
 #include <integer/popcount.h>
 
 #define __CLC_BODY "popcount.inc"

diff --git a/libclc/amdgcn/lib/math/fmax.cl b/libclc/amdgcn/lib/math/fmax.cl
@@ -1,6 +1,5 @@
 #include <clc/clc.h>
-
-#include "../../../generic/lib/clcmacro.h"
+#include <clc/clcmacro.h>
 
 _CLC_DEF _CLC_OVERLOAD float fmax(float x, float y)
 {

diff --git a/libclc/amdgcn/lib/math/fmin.cl b/libclc/amdgcn/lib/math/fmin.cl
@@ -1,6 +1,5 @@
 #include <clc/clc.h>
-
-#include "../../../generic/lib/clcmacro.h"
+#include <clc/clcmacro.h>
 
 _CLC_DEF _CLC_OVERLOAD float fmin(float x, float y)
 {

diff --git a/libclc/amdgcn/lib/math/ldexp.cl b/libclc/amdgcn/lib/math/ldexp.cl
@@ -21,8 +21,7 @@
  */
 
 #include <clc/clc.h>
-
-#include "../../../generic/lib/clcmacro.h"
+#include <clc/clcmacro.h>
 
 #ifdef __HAS_LDEXPF__
 #define BUILTINF __builtin_amdgcn_ldexpf

diff --git a/libclc/amdgpu/lib/math/half_native_unary.inc b/libclc/amdgpu/lib/math/half_native_unary.inc
@@ -1,4 +1,4 @@
-#include <utils.h>
+#include <clc/utils.h>
 
 #define __CLC_HALF_FUNC(x) __CLC_CONCAT(half_, x)
 #define __CLC_NATIVE_FUNC(x) __CLC_CONCAT(native_, x)

diff --git a/libclc/amdgpu/lib/math/nextafter.cl b/libclc/amdgpu/lib/math/nextafter.cl
@@ -1,5 +1,5 @@
 #include <clc/clc.h>
-#include "../lib/clcmacro.h"
+#include <clc/clcmacro.h>
 #include <math/clc_nextafter.h>
 
 _CLC_DEFINE_BINARY_BUILTIN(float, nextafter, __clc_nextafter, float, float)

diff --git a/libclc/amdgpu/lib/math/sqrt.cl b/libclc/amdgpu/lib/math/sqrt.cl
@@ -20,9 +20,9 @@
  * THE SOFTWARE.
  */
 
-#include <clc/clc.h>
-#include "../../../generic/lib/clcmacro.h"
 #include "math/clc_sqrt.h"
+#include <clc/clc.h>
+#include <clc/clcmacro.h>
 
 _CLC_DEFINE_UNARY_BUILTIN(float, sqrt, __clc_sqrt, float)
 

diff --git a/libclc/generic/lib/clcmacro.h → libclc/clc/include/clc/clcmacro.h b/libclc/generic/lib/clcmacro.h → libclc/clc/include/clc/clcmacro.h
@@ -1,92 +1,105 @@
-#include <clc/clc.h>
-#include <utils.h>
-
-#define _CLC_UNARY_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE) \
-  DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE##2 x) { \
-    return (RET_TYPE##2)(FUNCTION(x.x), FUNCTION(x.y)); \
-  } \
-\
-  DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x) { \
-    return (RET_TYPE##3)(FUNCTION(x.x), FUNCTION(x.y), FUNCTION(x.z)); \
-  } \
-\
-  DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x) { \
-    return (RET_TYPE##4)(FUNCTION(x.lo), FUNCTION(x.hi)); \
-  } \
-\
-  DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x) { \
-    return (RET_TYPE##8)(FUNCTION(x.lo), FUNCTION(x.hi)); \
-  } \
-\
-  DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x) { \
-    return (RET_TYPE##16)(FUNCTION(x.lo), FUNCTION(x.hi)); \
+#ifndef __CLC_CLCMACRO_H__
+#define __CLC_CLCMACRO_H__
+
+#include <clc/internal/clc.h>
+#include <clc/utils.h>
+
+#define _CLC_UNARY_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE)          \
+  DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE##2 x) {                              \
+    return (RET_TYPE##2)(FUNCTION(x.x), FUNCTION(x.y));                        \
+  }                                                                            \
+                                                                               \
+  DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x) {                              \
+    return (RET_TYPE##3)(FUNCTION(x.x), FUNCTION(x.y), FUNCTION(x.z));         \
+  }                                                                            \
+                                                                               \
+  DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x) {                              \
+    return (RET_TYPE##4)(FUNCTION(x.lo), FUNCTION(x.hi));                      \
+  }                                                                            \
+                                                                               \
+  DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x) {                              \
+    return (RET_TYPE##8)(FUNCTION(x.lo), FUNCTION(x.hi));                      \
+  }                                                                            \
+                                                                               \
+  DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x) {                            \
+    return (RET_TYPE##16)(FUNCTION(x.lo), FUNCTION(x.hi));                     \
   }
 
-#define _CLC_BINARY_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, ARG2_TYPE) \
-  DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE##2 x, ARG2_TYPE##2 y) { \
-    return (RET_TYPE##2)(FUNCTION(x.x, y.x), FUNCTION(x.y, y.y)); \
-  } \
-\
-  DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x, ARG2_TYPE##3 y) { \
-    return (RET_TYPE##3)(FUNCTION(x.x, y.x), FUNCTION(x.y, y.y), \
-                         FUNCTION(x.z, y.z)); \
-  } \
-\
-  DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x, ARG2_TYPE##4 y) { \
-    return (RET_TYPE##4)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi)); \
-  } \
-\
-  DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x, ARG2_TYPE##8 y) { \
-    return (RET_TYPE##8)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi)); \
-  } \
-\
-  DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x, ARG2_TYPE##16 y) { \
-    return (RET_TYPE##16)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi)); \
+#define _CLC_BINARY_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE,         \
+                              ARG2_TYPE)                                       \
+  DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE##2 x, ARG2_TYPE##2 y) {              \
+    return (RET_TYPE##2)(FUNCTION(x.x, y.x), FUNCTION(x.y, y.y));              \
+  }                                                                            \
+                                                                               \
+  DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x, ARG2_TYPE##3 y) {              \
+    return (RET_TYPE##3)(FUNCTION(x.x, y.x), FUNCTION(x.y, y.y),               \
+                         FUNCTION(x.z, y.z));                                  \
+  }                                                                            \
+                                                                               \
+  DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x, ARG2_TYPE##4 y) {              \
+    return (RET_TYPE##4)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi));          \
+  }                                                                            \
+                                                                               \
+  DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x, ARG2_TYPE##8 y) {              \
+    return (RET_TYPE##8)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi));          \
+  }                                                                            \
+                                                                               \
+  DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x, ARG2_TYPE##16 y) {           \
+    return (RET_TYPE##16)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi));         \
   }
 
-#define _CLC_V_S_V_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, ARG2_TYPE) \
-  DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE x, ARG2_TYPE##2 y) { \
-    return (RET_TYPE##2)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
-  } \
-\
-  DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE x, ARG2_TYPE##3 y) { \
-    return (RET_TYPE##3)(FUNCTION(x, y.x), FUNCTION(x, y.y), \
-                         FUNCTION(x, y.z)); \
-  } \
-\
-  DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE x, ARG2_TYPE##4 y) { \
-    return (RET_TYPE##4)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
-  } \
-\
-  DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE x, ARG2_TYPE##8 y) { \
-    return (RET_TYPE##8)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
-  } \
-\
-  DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE x, ARG2_TYPE##16 y) { \
-    return (RET_TYPE##16)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
-  } \
-\
-
-#define _CLC_TERNARY_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, ARG2_TYPE, ARG3_TYPE) \
-  DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE##2 x, ARG2_TYPE##2 y, ARG3_TYPE##2 z) { \
-    return (RET_TYPE##2)(FUNCTION(x.x, y.x, z.x), FUNCTION(x.y, y.y, z.y)); \
-  } \
-\
-  DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x, ARG2_TYPE##3 y, ARG3_TYPE##3 z) { \
-    return (RET_TYPE##3)(FUNCTION(x.x, y.x, z.x), FUNCTION(x.y, y.y, z.y), \
-                         FUNCTION(x.z, y.z, z.z)); \
-  } \
-\
-  DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x, ARG2_TYPE##4 y, ARG3_TYPE##4 z) { \
-    return (RET_TYPE##4)(FUNCTION(x.lo, y.lo, z.lo), FUNCTION(x.hi, y.hi, z.hi)); \
-  } \
-\
-  DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x, ARG2_TYPE##8 y, ARG3_TYPE##8 z) { \
-    return (RET_TYPE##8)(FUNCTION(x.lo, y.lo, z.lo), FUNCTION(x.hi, y.hi, z.hi)); \
-  } \
-\
-  DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x, ARG2_TYPE##16 y, ARG3_TYPE##16 z) { \
-    return (RET_TYPE##16)(FUNCTION(x.lo, y.lo, z.lo), FUNCTION(x.hi, y.hi, z.hi)); \
+#define _CLC_V_S_V_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE,          \
+                             ARG2_TYPE)                                        \
+  DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE x, ARG2_TYPE##2 y) {                 \
+    return (RET_TYPE##2)(FUNCTION(x, y.lo), FUNCTION(x, y.hi));                \
+  }                                                                            \
+                                                                               \
+  DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE x, ARG2_TYPE##3 y) {                 \
+    return (RET_TYPE##3)(FUNCTION(x, y.x), FUNCTION(x, y.y),                   \
+                         FUNCTION(x, y.z));                                    \
+  }                                                                            \
+                                                                               \
+  DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE x, ARG2_TYPE##4 y) {                 \
+    return (RET_TYPE##4)(FUNCTION(x, y.lo), FUNCTION(x, y.hi));                \
+  }                                                                            \
+                                                                               \
+  DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE x, ARG2_TYPE##8 y) {                 \
+    return (RET_TYPE##8)(FUNCTION(x, y.lo), FUNCTION(x, y.hi));                \
+  }                                                                            \
+                                                                               \
+  DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE x, ARG2_TYPE##16 y) {               \
+    return (RET_TYPE##16)(FUNCTION(x, y.lo), FUNCTION(x, y.hi));               \
+  }
+
+#define _CLC_TERNARY_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE,        \
+                               ARG2_TYPE, ARG3_TYPE)                           \
+  DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE##2 x, ARG2_TYPE##2 y,                \
+                                ARG3_TYPE##2 z) {                              \
+    return (RET_TYPE##2)(FUNCTION(x.x, y.x, z.x), FUNCTION(x.y, y.y, z.y));    \
+  }                                                                            \
+                                                                               \
+  DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x, ARG2_TYPE##3 y,                \
+                                ARG3_TYPE##3 z) {                              \
+    return (RET_TYPE##3)(FUNCTION(x.x, y.x, z.x), FUNCTION(x.y, y.y, z.y),     \
+                         FUNCTION(x.z, y.z, z.z));                             \
+  }                                                                            \
+                                                                               \
+  DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x, ARG2_TYPE##4 y,                \
+                                ARG3_TYPE##4 z) {                              \
+    return (RET_TYPE##4)(FUNCTION(x.lo, y.lo, z.lo),                           \
+                         FUNCTION(x.hi, y.hi, z.hi));                          \
+  }                                                                            \
+                                                                               \
+  DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x, ARG2_TYPE##8 y,                \
+                                ARG3_TYPE##8 z) {                              \
+    return (RET_TYPE##8)(FUNCTION(x.lo, y.lo, z.lo),                           \
+                         FUNCTION(x.hi, y.hi, z.hi));                          \
+  }                                                                            \
+                                                                               \
+  DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x, ARG2_TYPE##16 y,             \
+                                 ARG3_TYPE##16 z) {                            \
+    return (RET_TYPE##16)(FUNCTION(x.lo, y.lo, z.lo),                          \
+                          FUNCTION(x.hi, y.hi, z.hi));                         \
   }
 
 #define _CLC_V_S_S_V_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE,        \
@@ -161,21 +174,24 @@
                            ARG2_TYPE, 8) *)((ADDR_SPACE ARG2_TYPE *)y + 8)));  \
   }
 
-#define _CLC_DEFINE_BINARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE, ARG2_TYPE) \
-_CLC_DEF _CLC_OVERLOAD RET_TYPE FUNCTION(ARG1_TYPE x, ARG2_TYPE y) { \
-  return BUILTIN(x, y); \
-} \
-_CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE, ARG2_TYPE)
+#define _CLC_DEFINE_BINARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE,     \
+                                   ARG2_TYPE)                                  \
+  _CLC_DEF _CLC_OVERLOAD RET_TYPE FUNCTION(ARG1_TYPE x, ARG2_TYPE y) {         \
+    return BUILTIN(x, y);                                                      \
+  }                                                                            \
+  _CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE, \
+                        ARG2_TYPE)
 
-#define _CLC_DEFINE_BINARY_BUILTIN_WITH_SCALAR_SECOND_ARG(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE, ARG2_TYPE) \
-_CLC_DEFINE_BINARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE, ARG2_TYPE) \
-_CLC_BINARY_VECTORIZE_SCALAR_SECOND_ARG(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE, ARG2_TYPE)
+#define _CLC_DEFINE_BINARY_BUILTIN_WITH_SCALAR_SECOND_ARG(                     \
+    RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE, ARG2_TYPE)                         \
+  _CLC_DEFINE_BINARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE,           \
+                             ARG2_TYPE)                                        \
+  _CLC_BINARY_VECTORIZE_SCALAR_SECOND_ARG(_CLC_OVERLOAD _CLC_DEF, RET_TYPE,    \
+                                          FUNCTION, ARG1_TYPE, ARG2_TYPE)
 
-#define _CLC_DEFINE_UNARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE) \
-_CLC_DEF _CLC_OVERLOAD RET_TYPE FUNCTION(ARG1_TYPE x) { \
-  return BUILTIN(x); \
-} \
-_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE)
+#define _CLC_DEFINE_UNARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE)      \
+  _CLC_DEF _CLC_OVERLOAD RET_TYPE FUNCTION(ARG1_TYPE x) { return BUILTIN(x); } \
+  _CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE)
 
 #ifdef cl_khr_fp16
 
@@ -199,3 +215,5 @@ _CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE)
 #define _CLC_DEFINE_BINARY_BUILTIN_FP16(FUNCTION)
 
 #endif
+
+#endif // __CLC_CLCMACRO_H__
diff --git a/libclc/clspv/lib/math/fma.cl b/libclc/clspv/lib/math/fma.cl
@@ -24,9 +24,9 @@
 // (__clc_sw_fma), but avoids the use of ulong in favor of uint2. The logic has
 // been updated as appropriate.
 
-#include <clc/clc.h>
-#include "../../../generic/lib/clcmacro.h"
 #include "../../../generic/lib/math/math.h"
+#include <clc/clc.h>
+#include <clc/clcmacro.h>
 
 struct fp {
   uint2 mantissa;

diff --git a/libclc/generic/include/math/clc_sqrt.h b/libclc/generic/include/math/clc_sqrt.h
@@ -1,3 +1,6 @@
+#include <clc/clcfunc.h>
+#include <clc/clctypes.h>
+
 #define __CLC_FUNCTION __clc_sqrt
 #define __CLC_BODY <clc/math/unary_decl.inc>
 #include <clc/math/gentype.inc>

diff --git a/libclc/generic/include/utils.h b/libclc/generic/include/utils.h
diff --git a/libclc/generic/lib/atom_int32_binary.inc b/libclc/generic/lib/atom_int32_binary.inc
@@ -1,5 +1,5 @@
 #include <clc/clc.h>
-#include "utils.h"
+#include <clc/utils.h>
 
 #define __CLC_ATOM_IMPL(AS, TYPE) \
 _CLC_OVERLOAD _CLC_DEF TYPE __CLC_XCONCAT(atom_, __CLC_ATOMIC_OP) (volatile AS TYPE *p, TYPE val) { \

diff --git a/libclc/generic/lib/common/degrees.cl b/libclc/generic/lib/common/degrees.cl
@@ -21,8 +21,7 @@
  */
 
 #include <clc/clc.h>
-
-#include "../clcmacro.h"
+#include <clc/clcmacro.h>
 
 _CLC_OVERLOAD _CLC_DEF float degrees(float radians) {
   // 180/pi = ~57.29577951308232087685 or 0x1.ca5dc1a63c1f8p+5 or 0x1.ca5dc2p+5F

diff --git a/libclc/generic/lib/common/radians.cl b/libclc/generic/lib/common/radians.cl
@@ -21,8 +21,7 @@
  */
 
 #include <clc/clc.h>
-
-#include "../clcmacro.h"
+#include <clc/clcmacro.h>
 
 _CLC_OVERLOAD _CLC_DEF float radians(float degrees) {
   // pi/180 = ~0.01745329251994329577 or 0x1.1df46a2529d39p-6 or 0x1.1df46ap-6F

diff --git a/libclc/generic/lib/common/sign.cl b/libclc/generic/lib/common/sign.cl
@@ -1,5 +1,5 @@
 #include <clc/clc.h>
-#include "../clcmacro.h"
+#include <clc/clcmacro.h>
 
 #define SIGN(TYPE, F) \
 _CLC_DEF _CLC_OVERLOAD TYPE sign(TYPE x) { \

diff --git a/libclc/generic/lib/common/smoothstep.cl b/libclc/generic/lib/common/smoothstep.cl
@@ -21,8 +21,7 @@
  */
 
 #include <clc/clc.h>
-
-#include "../clcmacro.h"
+#include <clc/clcmacro.h>
 
 _CLC_OVERLOAD _CLC_DEF float smoothstep(float edge0, float edge1, float x) {
   float t = clamp((x - edge0) / (edge1 - edge0), 0.0f, 1.0f);

diff --git a/libclc/generic/lib/common/step.cl b/libclc/generic/lib/common/step.cl
@@ -21,8 +21,7 @@
  */
 
 #include <clc/clc.h>
-
-#include "../clcmacro.h"
+#include <clc/clcmacro.h>
 
 _CLC_OVERLOAD _CLC_DEF float step(float edge, float x) {
   return x < edge ? 0.0f : 1.0f;

diff --git a/libclc/generic/lib/integer/add_sat.cl b/libclc/generic/lib/integer/add_sat.cl
@@ -1,5 +1,5 @@
 #include <clc/clc.h>
-#include "../clcmacro.h"
+#include <clc/clcmacro.h>
 
 // From add_sat.ll
 _CLC_DECL char   __clc_add_sat_s8(char, char);

diff --git a/libclc/generic/lib/integer/clz.cl b/libclc/generic/lib/integer/clz.cl
@@ -1,5 +1,5 @@
 #include <clc/clc.h>
-#include "../clcmacro.h"
+#include <clc/clcmacro.h>
 
 _CLC_OVERLOAD _CLC_DEF char clz(char x) {
   return clz((ushort)(uchar)x) - 8;

diff --git a/libclc/generic/lib/integer/mad_sat.cl b/libclc/generic/lib/integer/mad_sat.cl
@@ -1,5 +1,5 @@
 #include <clc/clc.h>
-#include "../clcmacro.h"
+#include <clc/clcmacro.h>
 
 _CLC_OVERLOAD _CLC_DEF char mad_sat(char x, char y, char z) {
   return clamp((short)mad24((short)x, (short)y, (short)z), (short)CHAR_MIN, (short) CHAR_MAX);