[libclc] Add __attribute__((const)) to functions that don't access memory (#152456)
Before this PR, PostOrderFunctionAttrsPass in opt run can deduce
memory(none) for these functions.
This PR explicitly adds the attribute to align with Clang's OpenCL
headers and ensures the attribute is present throughout the compilation
flow. Generated bitcode files amdgcn--amdhsa.bc and nvptx64--nvidiacl.bc
become slightly smaller.
diff --git a/libclc/clc/include/clc/clcfunc.h b/libclc/clc/include/clc/clcfunc.h
index 10b9cdd..30feaf9 100644
--- a/libclc/clc/include/clc/clcfunc.h
+++ b/libclc/clc/include/clc/clcfunc.h
@@ -12,6 +12,7 @@
#define _CLC_OVERLOAD __attribute__((overloadable))
#define _CLC_DECL
#define _CLC_INLINE __attribute__((always_inline)) inline
+#define _CLC_CONST __attribute__((const))
// avoid inlines for SPIR-V related targets since we'll optimise later in the
// chain
diff --git a/libclc/clc/include/clc/common/clc_smoothstep.inc b/libclc/clc/include/clc/common/clc_smoothstep.inc
index 158c578..e3752e8 100644
--- a/libclc/clc/include/clc/common/clc_smoothstep.inc
+++ b/libclc/clc/include/clc/common/clc_smoothstep.inc
@@ -6,6 +6,5 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE __clc_smoothstep(__CLC_GENTYPE edge0,
- __CLC_GENTYPE edge1,
- __CLC_GENTYPE x);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE
+__clc_smoothstep(__CLC_GENTYPE edge0, __CLC_GENTYPE edge1, __CLC_GENTYPE x);
diff --git a/libclc/clc/include/clc/geometric/binary_decl.inc b/libclc/clc/include/clc/geometric/binary_decl.inc
index 342d97a..4a4235a 100644
--- a/libclc/clc/include/clc/geometric/binary_decl.inc
+++ b/libclc/clc/include/clc/geometric/binary_decl.inc
@@ -10,7 +10,7 @@
#if (__CLC_VECSIZE_OR_1 == 1 || __CLC_VECSIZE_OR_1 == 2 || \
__CLC_VECSIZE_OR_1 == 3 || __CLC_VECSIZE_OR_1 == 4)
-_CLC_OVERLOAD _CLC_DECL __CLC_SCALAR_GENTYPE FUNCTION(__CLC_GENTYPE a,
- __CLC_GENTYPE b);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_SCALAR_GENTYPE
+FUNCTION(__CLC_GENTYPE a, __CLC_GENTYPE b);
#endif
diff --git a/libclc/clc/include/clc/geometric/clc_cross.h b/libclc/clc/include/clc/geometric/clc_cross.h
index e5aa913..c0a9cac 100644
--- a/libclc/clc/include/clc/geometric/clc_cross.h
+++ b/libclc/clc/include/clc/geometric/clc_cross.h
@@ -11,22 +11,22 @@
#include <clc/internal/clc.h>
-_CLC_OVERLOAD _CLC_DECL float3 __clc_cross(float3 p0, float3 p1);
-_CLC_OVERLOAD _CLC_DECL float4 __clc_cross(float4 p0, float4 p1);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL float3 __clc_cross(float3 p0, float3 p1);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL float4 __clc_cross(float4 p0, float4 p1);
#ifdef cl_khr_fp64
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
-_CLC_OVERLOAD _CLC_DECL double3 __clc_cross(double3 p0, double3 p1);
-_CLC_OVERLOAD _CLC_DECL double4 __clc_cross(double4 p0, double4 p1);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL double3 __clc_cross(double3 p0, double3 p1);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL double4 __clc_cross(double4 p0, double4 p1);
#endif
#ifdef cl_khr_fp16
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
-_CLC_OVERLOAD _CLC_DECL half3 __clc_cross(half3 p0, half3 p1);
-_CLC_OVERLOAD _CLC_DECL half4 __clc_cross(half4 p0, half4 p1);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL half3 __clc_cross(half3 p0, half3 p1);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL half4 __clc_cross(half4 p0, half4 p1);
#endif
diff --git a/libclc/clc/include/clc/geometric/unary_decl.inc b/libclc/clc/include/clc/geometric/unary_decl.inc
index cff6600..bca1f3b 100644
--- a/libclc/clc/include/clc/geometric/unary_decl.inc
+++ b/libclc/clc/include/clc/geometric/unary_decl.inc
@@ -10,7 +10,7 @@
#if (__CLC_VECSIZE_OR_1 == 1 || __CLC_VECSIZE_OR_1 == 2 || \
__CLC_VECSIZE_OR_1 == 3 || __CLC_VECSIZE_OR_1 == 4)
-_CLC_OVERLOAD _CLC_DECL
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL
#ifdef __CLC_GEOMETRIC_RET_GENTYPE
__CLC_GENTYPE
#else
diff --git a/libclc/clc/include/clc/integer/clc_abs.inc b/libclc/clc/include/clc/integer/clc_abs.inc
index f1e2c5a..4c34b0d 100644
--- a/libclc/clc/include/clc/integer/clc_abs.inc
+++ b/libclc/clc/include/clc/integer/clc_abs.inc
@@ -6,4 +6,4 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_U_GENTYPE __clc_abs(__CLC_GENTYPE x);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_U_GENTYPE __clc_abs(__CLC_GENTYPE x);
diff --git a/libclc/clc/include/clc/integer/clc_abs_diff.inc b/libclc/clc/include/clc/integer/clc_abs_diff.inc
index 2514a37..84ef15c 100644
--- a/libclc/clc/include/clc/integer/clc_abs_diff.inc
+++ b/libclc/clc/include/clc/integer/clc_abs_diff.inc
@@ -6,5 +6,5 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_U_GENTYPE __clc_abs_diff(__CLC_GENTYPE x,
- __CLC_GENTYPE y);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_U_GENTYPE
+__clc_abs_diff(__CLC_GENTYPE x, __CLC_GENTYPE y);
diff --git a/libclc/clc/include/clc/integer/clc_bitfield_extract_decl.inc b/libclc/clc/include/clc/integer/clc_bitfield_extract_decl.inc
index c93eff0..b3f0e71 100644
--- a/libclc/clc/include/clc/integer/clc_bitfield_extract_decl.inc
+++ b/libclc/clc/include/clc/integer/clc_bitfield_extract_decl.inc
@@ -6,5 +6,5 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __RETTYPE FUNCTION(__CLC_GENTYPE base, uint offset,
- uint count);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __RETTYPE FUNCTION(__CLC_GENTYPE base,
+ uint offset, uint count);
diff --git a/libclc/clc/include/clc/integer/clc_bitfield_insert.inc b/libclc/clc/include/clc/integer/clc_bitfield_insert.inc
index 22f58bd..de0b3d9 100644
--- a/libclc/clc/include/clc/integer/clc_bitfield_insert.inc
+++ b/libclc/clc/include/clc/integer/clc_bitfield_insert.inc
@@ -6,6 +6,7 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_GENTYPE base,
- __CLC_GENTYPE insert,
- uint offset, uint count);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_GENTYPE base,
+ __CLC_GENTYPE insert,
+ uint offset,
+ uint count);
diff --git a/libclc/clc/include/clc/math/binary_decl_with_scalar_second_arg.inc b/libclc/clc/include/clc/math/binary_decl_with_scalar_second_arg.inc
index 1b2b9ff..45f39b7 100644
--- a/libclc/clc/include/clc/math/binary_decl_with_scalar_second_arg.inc
+++ b/libclc/clc/include/clc/math/binary_decl_with_scalar_second_arg.inc
@@ -6,7 +6,7 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_GENTYPE a,
- __CLC_GENTYPE b);
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_GENTYPE a,
- __CLC_SCALAR_GENTYPE b);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_GENTYPE a,
+ __CLC_GENTYPE b);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE
+FUNCTION(__CLC_GENTYPE a, __CLC_SCALAR_GENTYPE b);
diff --git a/libclc/clc/include/clc/math/binary_def_via_fp32.inc b/libclc/clc/include/clc/math/binary_def_via_fp32.inc
index b53a178..c8cdb1e 100644
--- a/libclc/clc/include/clc/math/binary_def_via_fp32.inc
+++ b/libclc/clc/include/clc/math/binary_def_via_fp32.inc
@@ -6,8 +6,8 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE FUNCTION(__CLC_GENTYPE x,
- __CLC_GENTYPE y) {
+_CLC_OVERLOAD _CLC_CONST _CLC_DEF __CLC_GENTYPE FUNCTION(__CLC_GENTYPE x,
+ __CLC_GENTYPE y) {
return __CLC_CONVERT_GENTYPE(
FUNCTION(__CLC_CONVERT_FLOATN(x), __CLC_CONVERT_FLOATN(y)));
}
diff --git a/libclc/clc/include/clc/math/clc_exp_helper.inc b/libclc/clc/include/clc/math/clc_exp_helper.inc
index cdf6504..cb2a0d3 100644
--- a/libclc/clc/include/clc/math/clc_exp_helper.inc
+++ b/libclc/clc/include/clc/math/clc_exp_helper.inc
@@ -6,8 +6,6 @@
//
//===----------------------------------------------------------------------===//
-_CLC_DECL _CLC_OVERLOAD __CLC_GENTYPE __clc_exp_helper(__CLC_GENTYPE x,
- __CLC_GENTYPE x_min,
- __CLC_GENTYPE x_max,
- __CLC_GENTYPE r,
- __CLC_INTN n);
+_CLC_DECL _CLC_OVERLOAD _CLC_CONST __CLC_GENTYPE
+__clc_exp_helper(__CLC_GENTYPE x, __CLC_GENTYPE x_min, __CLC_GENTYPE x_max,
+ __CLC_GENTYPE r, __CLC_INTN n);
diff --git a/libclc/clc/include/clc/math/clc_ldexp.inc b/libclc/clc/include/clc/math/clc_ldexp.inc
index 9b88359..b0bea29 100644
--- a/libclc/clc/include/clc/math/clc_ldexp.inc
+++ b/libclc/clc/include/clc/math/clc_ldexp.inc
@@ -6,4 +6,5 @@
//
//===----------------------------------------------------------------------===//
-_CLC_DECL _CLC_OVERLOAD __CLC_GENTYPE __clc_ldexp(__CLC_GENTYPE, __CLC_INTN);
+_CLC_DECL _CLC_OVERLOAD _CLC_CONST __CLC_GENTYPE __clc_ldexp(__CLC_GENTYPE,
+ __CLC_INTN);
diff --git a/libclc/clc/include/clc/math/clc_nan.inc b/libclc/clc/include/clc/math/clc_nan.inc
index 22604e0..9203499 100644
--- a/libclc/clc/include/clc/math/clc_nan.inc
+++ b/libclc/clc/include/clc/math/clc_nan.inc
@@ -6,5 +6,5 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __clc_nan(__CLC_U_GENTYPE code);
-_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __clc_nan(__CLC_S_GENTYPE code);
+_CLC_OVERLOAD _CLC_CONST _CLC_DEF __CLC_GENTYPE __clc_nan(__CLC_U_GENTYPE code);
+_CLC_OVERLOAD _CLC_CONST _CLC_DEF __CLC_GENTYPE __clc_nan(__CLC_S_GENTYPE code);
diff --git a/libclc/clc/include/clc/math/unary_decl.inc b/libclc/clc/include/clc/math/unary_decl.inc
index ce1a71b..46108c5 100644
--- a/libclc/clc/include/clc/math/unary_decl.inc
+++ b/libclc/clc/include/clc/math/unary_decl.inc
@@ -6,4 +6,4 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_GENTYPE x);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_GENTYPE x);
diff --git a/libclc/clc/include/clc/math/unary_decl_with_int_return.inc b/libclc/clc/include/clc/math/unary_decl_with_int_return.inc
index ed8d0a1..64fcf42 100644
--- a/libclc/clc/include/clc/math/unary_decl_with_int_return.inc
+++ b/libclc/clc/include/clc/math/unary_decl_with_int_return.inc
@@ -6,4 +6,4 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_INTN FUNCTION(__CLC_GENTYPE x);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_INTN FUNCTION(__CLC_GENTYPE x);
diff --git a/libclc/clc/include/clc/misc/shuffle2_decl.inc b/libclc/clc/include/clc/misc/shuffle2_decl.inc
index 47fd5e2..3504e78 100644
--- a/libclc/clc/include/clc/misc/shuffle2_decl.inc
+++ b/libclc/clc/include/clc/misc/shuffle2_decl.inc
@@ -12,16 +12,16 @@
// The return type is same base type as the input type, with the same vector
// size as the mask. Elements in the mask must be the same size (number of bits)
// as the input value., e.g. char8 ret = shuffle(char2 x, uchar8 mask);
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE
FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 2) x,
__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 2) y, __CLC_U_GENTYPE mask);
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE
FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 4) x,
__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 4) y, __CLC_U_GENTYPE mask);
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE
FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 8) x,
__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 8) y, __CLC_U_GENTYPE mask);
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE
FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 16) x,
__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 16) y, __CLC_U_GENTYPE mask);
diff --git a/libclc/clc/include/clc/misc/shuffle_decl.inc b/libclc/clc/include/clc/misc/shuffle_decl.inc
index c6c64d5..1445aaf 100644
--- a/libclc/clc/include/clc/misc/shuffle_decl.inc
+++ b/libclc/clc/include/clc/misc/shuffle_decl.inc
@@ -12,13 +12,13 @@
// The return type is same base type as the input type, with the same vector
// size as the mask. Elements in the mask must be the same size (number of bits)
// as the input value., e.g. char8 ret = shuffle(char2 x, uchar8 mask);
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE
FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 2) x, __CLC_U_GENTYPE mask);
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE
FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 4) x, __CLC_U_GENTYPE mask);
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE
FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 8) x, __CLC_U_GENTYPE mask);
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE
FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 16) x, __CLC_U_GENTYPE mask);
#endif
diff --git a/libclc/clc/include/clc/relational/binary_decl.inc b/libclc/clc/include/clc/relational/binary_decl.inc
index dc8ec9d..87cbc8c 100644
--- a/libclc/clc/include/clc/relational/binary_decl.inc
+++ b/libclc/clc/include/clc/relational/binary_decl.inc
@@ -12,6 +12,7 @@
#define __RETTYPE __CLC_BIT_INTN
#endif
-_CLC_OVERLOAD _CLC_DECL __RETTYPE FUNCTION(__CLC_GENTYPE a, __CLC_GENTYPE b);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __RETTYPE FUNCTION(__CLC_GENTYPE a,
+ __CLC_GENTYPE b);
#undef __RETTYPE
diff --git a/libclc/clc/include/clc/relational/clc_all.h b/libclc/clc/include/clc/relational/clc_all.h
index 4d01fb7..272d879 100644
--- a/libclc/clc/include/clc/relational/clc_all.h
+++ b/libclc/clc/include/clc/relational/clc_all.h
@@ -12,7 +12,8 @@
#include <clc/clcfunc.h>
#include <clc/clctypes.h>
-#define _CLC_ALL_DECL(TYPE) _CLC_OVERLOAD _CLC_DECL int __clc_all(TYPE v);
+#define _CLC_ALL_DECL(TYPE) \
+ _CLC_OVERLOAD _CLC_CONST _CLC_DECL int __clc_all(TYPE v);
#define _CLC_VECTOR_ALL_DECL(TYPE) \
_CLC_ALL_DECL(TYPE) \
diff --git a/libclc/clc/include/clc/relational/clc_any.h b/libclc/clc/include/clc/relational/clc_any.h
index 1e287af..82d08c0 100644
--- a/libclc/clc/include/clc/relational/clc_any.h
+++ b/libclc/clc/include/clc/relational/clc_any.h
@@ -12,7 +12,8 @@
#include <clc/clcfunc.h>
#include <clc/clctypes.h>
-#define _CLC_ANY_DECL(TYPE) _CLC_OVERLOAD _CLC_DECL int __clc_any(TYPE v);
+#define _CLC_ANY_DECL(TYPE) \
+ _CLC_OVERLOAD _CLC_CONST _CLC_DECL int __clc_any(TYPE v);
#define _CLC_VECTOR_ANY_DECL(TYPE) \
_CLC_ANY_DECL(TYPE) \
diff --git a/libclc/clc/include/clc/relational/clc_bitselect.inc b/libclc/clc/include/clc/relational/clc_bitselect.inc
index b81cfc2..564f3b0 100644
--- a/libclc/clc/include/clc/relational/clc_bitselect.inc
+++ b/libclc/clc/include/clc/relational/clc_bitselect.inc
@@ -6,6 +6,5 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE __clc_bitselect(__CLC_GENTYPE x,
- __CLC_GENTYPE y,
- __CLC_GENTYPE z);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE
+__clc_bitselect(__CLC_GENTYPE x, __CLC_GENTYPE y, __CLC_GENTYPE z);
diff --git a/libclc/clc/include/clc/relational/clc_isequal.h b/libclc/clc/include/clc/relational/clc_isequal.h
index f3789f5..2585dbd 100644
--- a/libclc/clc/include/clc/relational/clc_isequal.h
+++ b/libclc/clc/include/clc/relational/clc_isequal.h
@@ -13,7 +13,7 @@
#include <clc/clctypes.h>
#define _CLC_ISEQUAL_DECL(TYPE, RETTYPE) \
- _CLC_OVERLOAD _CLC_DECL RETTYPE __clc_isequal(TYPE x, TYPE y);
+ _CLC_OVERLOAD _CLC_CONST _CLC_DECL RETTYPE __clc_isequal(TYPE x, TYPE y);
#define _CLC_VECTOR_ISEQUAL_DECL(TYPE, RETTYPE) \
_CLC_ISEQUAL_DECL(TYPE##2, RETTYPE##2) \
diff --git a/libclc/clc/include/clc/relational/clc_isinf.h b/libclc/clc/include/clc/relational/clc_isinf.h
index 39a9de6..fcd87e94 100644
--- a/libclc/clc/include/clc/relational/clc_isinf.h
+++ b/libclc/clc/include/clc/relational/clc_isinf.h
@@ -13,7 +13,7 @@
#include <clc/clctypes.h>
#define _CLC_ISINF_DECL(RET_TYPE, ARG_TYPE) \
- _CLC_OVERLOAD _CLC_DECL RET_TYPE __clc_isinf(ARG_TYPE);
+ _CLC_OVERLOAD _CLC_CONST _CLC_DECL RET_TYPE __clc_isinf(ARG_TYPE);
#define _CLC_VECTOR_ISINF_DECL(RET_TYPE, ARG_TYPE) \
_CLC_ISINF_DECL(RET_TYPE##2, ARG_TYPE##2) \
diff --git a/libclc/clc/include/clc/relational/clc_isnan.h b/libclc/clc/include/clc/relational/clc_isnan.h
index d0821a3..779e6cc 100644
--- a/libclc/clc/include/clc/relational/clc_isnan.h
+++ b/libclc/clc/include/clc/relational/clc_isnan.h
@@ -13,7 +13,7 @@
#include <clc/clctypes.h>
#define _CLC_ISNAN_DECL(RET_TYPE, ARG_TYPE) \
- _CLC_OVERLOAD _CLC_DECL RET_TYPE __clc_isnan(ARG_TYPE);
+ _CLC_OVERLOAD _CLC_CONST _CLC_DECL RET_TYPE __clc_isnan(ARG_TYPE);
#define _CLC_VECTOR_ISNAN_DECL(RET_TYPE, ARG_TYPE) \
_CLC_ISNAN_DECL(RET_TYPE##2, ARG_TYPE##2) \
diff --git a/libclc/clc/include/clc/relational/clc_select_decl.inc b/libclc/clc/include/clc/relational/clc_select_decl.inc
index 1cf026d..a2350f7 100644
--- a/libclc/clc/include/clc/relational/clc_select_decl.inc
+++ b/libclc/clc/include/clc/relational/clc_select_decl.inc
@@ -6,9 +6,7 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE __CLC_SELECT_FN(__CLC_GENTYPE x,
- __CLC_GENTYPE y,
- __CLC_S_GENTYPE z);
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE __CLC_SELECT_FN(__CLC_GENTYPE x,
- __CLC_GENTYPE y,
- __CLC_U_GENTYPE z);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE
+__CLC_SELECT_FN(__CLC_GENTYPE x, __CLC_GENTYPE y, __CLC_S_GENTYPE z);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE
+__CLC_SELECT_FN(__CLC_GENTYPE x, __CLC_GENTYPE y, __CLC_U_GENTYPE z);
diff --git a/libclc/clc/include/clc/relational/unary_decl.inc b/libclc/clc/include/clc/relational/unary_decl.inc
index cc3f2d0..f8123ee 100644
--- a/libclc/clc/include/clc/relational/unary_decl.inc
+++ b/libclc/clc/include/clc/relational/unary_decl.inc
@@ -12,6 +12,6 @@
#define __RETTYPE __CLC_BIT_INTN
#endif
-_CLC_OVERLOAD _CLC_DECL __RETTYPE FUNCTION(__CLC_GENTYPE x);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __RETTYPE FUNCTION(__CLC_GENTYPE x);
#undef __RETTYPE
diff --git a/libclc/clc/include/clc/shared/binary_decl.inc b/libclc/clc/include/clc/shared/binary_decl.inc
index ff4739d..c6e2b2a 100644
--- a/libclc/clc/include/clc/shared/binary_decl.inc
+++ b/libclc/clc/include/clc/shared/binary_decl.inc
@@ -6,5 +6,5 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_GENTYPE x,
- __CLC_GENTYPE y);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_GENTYPE x,
+ __CLC_GENTYPE y);
diff --git a/libclc/clc/include/clc/shared/binary_decl_with_int_second_arg.inc b/libclc/clc/include/clc/shared/binary_decl_with_int_second_arg.inc
index a86d89b..a091eb8 100644
--- a/libclc/clc/include/clc/shared/binary_decl_with_int_second_arg.inc
+++ b/libclc/clc/include/clc/shared/binary_decl_with_int_second_arg.inc
@@ -8,4 +8,5 @@
#include <clc/utils.h>
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_GENTYPE x, __CLC_INTN y);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_GENTYPE x,
+ __CLC_INTN y);
diff --git a/libclc/clc/include/clc/shared/clc_clamp.inc b/libclc/clc/include/clc/shared/clc_clamp.inc
index c1280a6..343fa9d 100644
--- a/libclc/clc/include/clc/shared/clc_clamp.inc
+++ b/libclc/clc/include/clc/shared/clc_clamp.inc
@@ -6,12 +6,11 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE __clc_clamp(__CLC_GENTYPE x,
- __CLC_GENTYPE y,
- __CLC_GENTYPE z);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE __clc_clamp(__CLC_GENTYPE x,
+ __CLC_GENTYPE y,
+ __CLC_GENTYPE z);
#ifndef __CLC_SCALAR
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE __clc_clamp(__CLC_GENTYPE x,
- __CLC_SCALAR_GENTYPE y,
- __CLC_SCALAR_GENTYPE z);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE
+__clc_clamp(__CLC_GENTYPE x, __CLC_SCALAR_GENTYPE y, __CLC_SCALAR_GENTYPE z);
#endif
diff --git a/libclc/clc/include/clc/shared/clc_max.inc b/libclc/clc/include/clc/shared/clc_max.inc
index ca99c12..e2a6de6 100644
--- a/libclc/clc/include/clc/shared/clc_max.inc
+++ b/libclc/clc/include/clc/shared/clc_max.inc
@@ -6,10 +6,10 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE __clc_max(__CLC_GENTYPE a,
- __CLC_GENTYPE b);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE __clc_max(__CLC_GENTYPE a,
+ __CLC_GENTYPE b);
#ifndef __CLC_SCALAR
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE __clc_max(__CLC_GENTYPE a,
- __CLC_SCALAR_GENTYPE b);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE
+__clc_max(__CLC_GENTYPE a, __CLC_SCALAR_GENTYPE b);
#endif
diff --git a/libclc/clc/include/clc/shared/clc_min.inc b/libclc/clc/include/clc/shared/clc_min.inc
index c7cfe23..86b5ede 100644
--- a/libclc/clc/include/clc/shared/clc_min.inc
+++ b/libclc/clc/include/clc/shared/clc_min.inc
@@ -6,10 +6,10 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE __clc_min(__CLC_GENTYPE a,
- __CLC_GENTYPE b);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE __clc_min(__CLC_GENTYPE a,
+ __CLC_GENTYPE b);
#ifndef __CLC_SCALAR
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE __clc_min(__CLC_GENTYPE a,
- __CLC_SCALAR_GENTYPE b);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE
+__clc_min(__CLC_GENTYPE a, __CLC_SCALAR_GENTYPE b);
#endif
diff --git a/libclc/clc/include/clc/shared/ternary_decl.inc b/libclc/clc/include/clc/shared/ternary_decl.inc
index a76db1e..1dc8672 100644
--- a/libclc/clc/include/clc/shared/ternary_decl.inc
+++ b/libclc/clc/include/clc/shared/ternary_decl.inc
@@ -6,5 +6,6 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_GENTYPE a, __CLC_GENTYPE b,
- __CLC_GENTYPE c);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_GENTYPE a,
+ __CLC_GENTYPE b,
+ __CLC_GENTYPE c);
diff --git a/libclc/clc/include/clc/shared/unary_decl.inc b/libclc/clc/include/clc/shared/unary_decl.inc
index ce1a71b..46108c5 100644
--- a/libclc/clc/include/clc/shared/unary_decl.inc
+++ b/libclc/clc/include/clc/shared/unary_decl.inc
@@ -6,4 +6,4 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_GENTYPE x);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_GENTYPE x);
diff --git a/libclc/clc/include/clc/workitem/clc_get_global_id.h b/libclc/clc/include/clc/workitem/clc_get_global_id.h
index da15ed8..8a78c78 100644
--- a/libclc/clc/include/clc/workitem/clc_get_global_id.h
+++ b/libclc/clc/include/clc/workitem/clc_get_global_id.h
@@ -11,6 +11,6 @@
#include <clc/internal/clc.h>
-_CLC_OVERLOAD _CLC_DECL size_t __clc_get_global_id(uint dim);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL size_t __clc_get_global_id(uint dim);
#endif // __CLC_WORKITEM_CLC_GET_GLOBAL_ID_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_global_offset.h b/libclc/clc/include/clc/workitem/clc_get_global_offset.h
index 639be48..bf6d4e5 100644
--- a/libclc/clc/include/clc/workitem/clc_get_global_offset.h
+++ b/libclc/clc/include/clc/workitem/clc_get_global_offset.h
@@ -11,6 +11,6 @@
#include <clc/internal/clc.h>
-_CLC_OVERLOAD _CLC_DECL size_t __clc_get_global_offset(uint dim);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL size_t __clc_get_global_offset(uint dim);
#endif // __CLC_WORKITEM_CLC_GET_GLOBAL_OFFSET_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_global_size.h b/libclc/clc/include/clc/workitem/clc_get_global_size.h
index fce11f4..87b88d7 100644
--- a/libclc/clc/include/clc/workitem/clc_get_global_size.h
+++ b/libclc/clc/include/clc/workitem/clc_get_global_size.h
@@ -11,6 +11,6 @@
#include <clc/internal/clc.h>
-_CLC_OVERLOAD _CLC_DECL size_t __clc_get_global_size(uint dim);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL size_t __clc_get_global_size(uint dim);
#endif // __CLC_WORKITEM_CLC_GET_GLOBAL_SIZE_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_group_id.h b/libclc/clc/include/clc/workitem/clc_get_group_id.h
index 3fd43cb..cc1b0f13 100644
--- a/libclc/clc/include/clc/workitem/clc_get_group_id.h
+++ b/libclc/clc/include/clc/workitem/clc_get_group_id.h
@@ -11,6 +11,6 @@
#include <clc/internal/clc.h>
-_CLC_OVERLOAD _CLC_DECL size_t __clc_get_group_id(uint dim);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL size_t __clc_get_group_id(uint dim);
#endif // __CLC_WORKITEM_CLC_GET_GROUP_ID_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_local_id.h b/libclc/clc/include/clc/workitem/clc_get_local_id.h
index c446127..a25e5eb 100644
--- a/libclc/clc/include/clc/workitem/clc_get_local_id.h
+++ b/libclc/clc/include/clc/workitem/clc_get_local_id.h
@@ -11,6 +11,6 @@
#include <clc/internal/clc.h>
-_CLC_OVERLOAD _CLC_DECL size_t __clc_get_local_id(uint dim);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL size_t __clc_get_local_id(uint dim);
#endif // __CLC_WORKITEM_CLC_GET_LOCAL_ID_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_local_linear_id.h b/libclc/clc/include/clc/workitem/clc_get_local_linear_id.h
index 34c47f2..9548644 100644
--- a/libclc/clc/include/clc/workitem/clc_get_local_linear_id.h
+++ b/libclc/clc/include/clc/workitem/clc_get_local_linear_id.h
@@ -11,6 +11,6 @@
#include <clc/internal/clc.h>
-_CLC_OVERLOAD _CLC_DECL size_t __clc_get_local_linear_id();
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL size_t __clc_get_local_linear_id();
#endif // __CLC_WORKITEM_CLC_GET_LOCAL_LINEAR_ID_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_local_size.h b/libclc/clc/include/clc/workitem/clc_get_local_size.h
index 7fe3c19..1272227 100644
--- a/libclc/clc/include/clc/workitem/clc_get_local_size.h
+++ b/libclc/clc/include/clc/workitem/clc_get_local_size.h
@@ -11,6 +11,6 @@
#include <clc/internal/clc.h>
-_CLC_OVERLOAD _CLC_DECL size_t __clc_get_local_size(uint dim);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL size_t __clc_get_local_size(uint dim);
#endif // __CLC_WORKITEM_CLC_GET_LOCAL_SIZE_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_max_sub_group_size.h b/libclc/clc/include/clc/workitem/clc_get_max_sub_group_size.h
index 8b521fe..c8ea2a1 100644
--- a/libclc/clc/include/clc/workitem/clc_get_max_sub_group_size.h
+++ b/libclc/clc/include/clc/workitem/clc_get_max_sub_group_size.h
@@ -11,6 +11,6 @@
#include <clc/internal/clc.h>
-_CLC_DEF _CLC_OVERLOAD uint __clc_get_max_sub_group_size();
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL uint __clc_get_max_sub_group_size();
#endif // __CLC_WORKITEM_CLC_GET_MAX_SUB_GROUP_SIZE_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_num_groups.h b/libclc/clc/include/clc/workitem/clc_get_num_groups.h
index f860944..6071fc3 100644
--- a/libclc/clc/include/clc/workitem/clc_get_num_groups.h
+++ b/libclc/clc/include/clc/workitem/clc_get_num_groups.h
@@ -11,6 +11,6 @@
#include <clc/internal/clc.h>
-_CLC_OVERLOAD _CLC_DECL size_t __clc_get_num_groups(uint dim);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL size_t __clc_get_num_groups(uint dim);
#endif // __CLC_WORKITEM_CLC_GET_NUM_GROUPS_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_num_sub_groups.h b/libclc/clc/include/clc/workitem/clc_get_num_sub_groups.h
index 6965aef..f94e898 100644
--- a/libclc/clc/include/clc/workitem/clc_get_num_sub_groups.h
+++ b/libclc/clc/include/clc/workitem/clc_get_num_sub_groups.h
@@ -11,6 +11,6 @@
#include <clc/internal/clc.h>
-_CLC_DEF _CLC_OVERLOAD uint __clc_get_num_sub_groups();
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL uint __clc_get_num_sub_groups();
#endif // __CLC_WORKITEM_CLC_GET_NUM_SUB_GROUPS_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_sub_group_id.h b/libclc/clc/include/clc/workitem/clc_get_sub_group_id.h
index ac3d7bd..61b54ec 100644
--- a/libclc/clc/include/clc/workitem/clc_get_sub_group_id.h
+++ b/libclc/clc/include/clc/workitem/clc_get_sub_group_id.h
@@ -11,6 +11,6 @@
#include <clc/internal/clc.h>
-_CLC_DEF _CLC_OVERLOAD uint __clc_get_sub_group_id();
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL uint __clc_get_sub_group_id();
#endif // __CLC_WORKITEM_CLC_GET_SUB_GROUP_ID_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_sub_group_local_id.h b/libclc/clc/include/clc/workitem/clc_get_sub_group_local_id.h
index 06bb6f8..a158153 100644
--- a/libclc/clc/include/clc/workitem/clc_get_sub_group_local_id.h
+++ b/libclc/clc/include/clc/workitem/clc_get_sub_group_local_id.h
@@ -11,6 +11,6 @@
#include <clc/internal/clc.h>
-_CLC_DEF _CLC_OVERLOAD uint __clc_get_sub_group_local_id();
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL uint __clc_get_sub_group_local_id();
#endif // __CLC_WORKITEM_CLC_GET_SUB_GROUP_LOCAL_ID_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_sub_group_size.h b/libclc/clc/include/clc/workitem/clc_get_sub_group_size.h
index a6e8e49..c4b1d92 100644
--- a/libclc/clc/include/clc/workitem/clc_get_sub_group_size.h
+++ b/libclc/clc/include/clc/workitem/clc_get_sub_group_size.h
@@ -11,6 +11,6 @@
#include <clc/internal/clc.h>
-_CLC_DEF _CLC_OVERLOAD uint __clc_get_sub_group_size();
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL uint __clc_get_sub_group_size();
#endif // __CLC_WORKITEM_CLC_GET_SUB_GROUP_SIZE_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_work_dim.h b/libclc/clc/include/clc/workitem/clc_get_work_dim.h
index 7e2da45..e93d377 100644
--- a/libclc/clc/include/clc/workitem/clc_get_work_dim.h
+++ b/libclc/clc/include/clc/workitem/clc_get_work_dim.h
@@ -11,6 +11,6 @@
#include <clc/internal/clc.h>
-_CLC_OVERLOAD _CLC_DECL uint __clc_get_work_dim();
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL uint __clc_get_work_dim();
#endif // __CLC_WORKITEM_CLC_GET_WORK_DIM_H__
diff --git a/libclc/opencl/include/clc/opencl/common/mix.inc b/libclc/opencl/include/clc/opencl/common/mix.inc
index 2b8350f..cd32085 100644
--- a/libclc/opencl/include/clc/opencl/common/mix.inc
+++ b/libclc/opencl/include/clc/opencl/common/mix.inc
@@ -6,10 +6,12 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE mix(__CLC_GENTYPE a, __CLC_GENTYPE b,
- __CLC_GENTYPE c);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE mix(__CLC_GENTYPE a,
+ __CLC_GENTYPE b,
+ __CLC_GENTYPE c);
#ifndef __CLC_SCALAR
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE mix(__CLC_GENTYPE a, __CLC_GENTYPE b,
- __CLC_SCALAR_GENTYPE c);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE mix(__CLC_GENTYPE a,
+ __CLC_GENTYPE b,
+ __CLC_SCALAR_GENTYPE c);
#endif
diff --git a/libclc/opencl/include/clc/opencl/common/smoothstep.inc b/libclc/opencl/include/clc/opencl/common/smoothstep.inc
index 4547483..f1402e3 100644
--- a/libclc/opencl/include/clc/opencl/common/smoothstep.inc
+++ b/libclc/opencl/include/clc/opencl/common/smoothstep.inc
@@ -6,13 +6,15 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE smoothstep(__CLC_GENTYPE edge0,
- __CLC_GENTYPE edge1,
- __CLC_GENTYPE x);
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE smoothstep(float edge0, float edge1,
- __CLC_GENTYPE x);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE smoothstep(__CLC_GENTYPE edge0,
+ __CLC_GENTYPE edge1,
+ __CLC_GENTYPE x);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE smoothstep(float edge0,
+ float edge1,
+ __CLC_GENTYPE x);
#ifdef cl_khr_fp64
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE smoothstep(double edge0, double edge1,
- __CLC_GENTYPE x);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE smoothstep(double edge0,
+ double edge1,
+ __CLC_GENTYPE x);
#endif
diff --git a/libclc/opencl/include/clc/opencl/common/step.inc b/libclc/opencl/include/clc/opencl/common/step.inc
index 86d4ed0..ab4b61e 100644
--- a/libclc/opencl/include/clc/opencl/common/step.inc
+++ b/libclc/opencl/include/clc/opencl/common/step.inc
@@ -6,9 +6,10 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE step(__CLC_GENTYPE edge, __CLC_GENTYPE x);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE step(__CLC_GENTYPE edge,
+ __CLC_GENTYPE x);
#ifndef __CLC_SCALAR
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE step(__CLC_SCALAR_GENTYPE edge,
- __CLC_GENTYPE x);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE step(__CLC_SCALAR_GENTYPE edge,
+ __CLC_GENTYPE x);
#endif
diff --git a/libclc/opencl/include/clc/opencl/geometric/cross.h b/libclc/opencl/include/clc/opencl/geometric/cross.h
index be13796..10e8f56 100644
--- a/libclc/opencl/include/clc/opencl/geometric/cross.h
+++ b/libclc/opencl/include/clc/opencl/geometric/cross.h
@@ -9,12 +9,12 @@
#ifndef __CLC_OPENCL_GEOMETRIC_CROSS_H__
#define __CLC_OPENCL_GEOMETRIC_CROSS_H__
-_CLC_OVERLOAD _CLC_DECL float3 cross(float3 p0, float3 p1);
-_CLC_OVERLOAD _CLC_DECL float4 cross(float4 p0, float4 p1);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL float3 cross(float3 p0, float3 p1);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL float4 cross(float4 p0, float4 p1);
#ifdef cl_khr_fp64
-_CLC_OVERLOAD _CLC_DECL double3 cross(double3 p0, double3 p1);
-_CLC_OVERLOAD _CLC_DECL double4 cross(double4 p0, double4 p1);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL double3 cross(double3 p0, double3 p1);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL double4 cross(double4 p0, double4 p1);
#endif
#endif // __CLC_OPENCL_GEOMETRIC_CROSS_H__
diff --git a/libclc/opencl/include/clc/opencl/integer/abs.inc b/libclc/opencl/include/clc/opencl/integer/abs.inc
index 6620ce3..352babb 100644
--- a/libclc/opencl/include/clc/opencl/integer/abs.inc
+++ b/libclc/opencl/include/clc/opencl/integer/abs.inc
@@ -6,4 +6,4 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_U_GENTYPE abs(__CLC_GENTYPE x);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_U_GENTYPE abs(__CLC_GENTYPE x);
diff --git a/libclc/opencl/include/clc/opencl/integer/abs_diff.inc b/libclc/opencl/include/clc/opencl/integer/abs_diff.inc
index 4187385..1f8fd63 100644
--- a/libclc/opencl/include/clc/opencl/integer/abs_diff.inc
+++ b/libclc/opencl/include/clc/opencl/integer/abs_diff.inc
@@ -6,5 +6,5 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_U_GENTYPE abs_diff(__CLC_GENTYPE x,
- __CLC_GENTYPE y);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_U_GENTYPE abs_diff(__CLC_GENTYPE x,
+ __CLC_GENTYPE y);
diff --git a/libclc/opencl/include/clc/opencl/integer/upsample.h b/libclc/opencl/include/clc/opencl/integer/upsample.h
index a713f2b..4432d49 100644
--- a/libclc/opencl/include/clc/opencl/integer/upsample.h
+++ b/libclc/opencl/include/clc/opencl/integer/upsample.h
@@ -12,7 +12,7 @@
#include <clc/opencl/opencl-base.h>
#define __CLC_UPSAMPLE_DECL(BGENTYPE, GENTYPE, UGENTYPE) \
- _CLC_OVERLOAD _CLC_DECL BGENTYPE upsample(GENTYPE hi, UGENTYPE lo);
+ _CLC_OVERLOAD _CLC_CONST _CLC_DECL BGENTYPE upsample(GENTYPE hi, UGENTYPE lo);
#define __CLC_UPSAMPLE_VEC(BGENTYPE, GENTYPE, UGENTYPE) \
__CLC_UPSAMPLE_DECL(BGENTYPE, GENTYPE, UGENTYPE) \
diff --git a/libclc/opencl/include/clc/opencl/math/ldexp.inc b/libclc/opencl/include/clc/opencl/math/ldexp.inc
index b5a5cfc..9f3ffeee 100644
--- a/libclc/opencl/include/clc/opencl/math/ldexp.inc
+++ b/libclc/opencl/include/clc/opencl/math/ldexp.inc
@@ -8,6 +8,6 @@
#ifndef __CLC_SCALAR
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE ldexp(__CLC_GENTYPE x, int n);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE ldexp(__CLC_GENTYPE x, int n);
#endif
diff --git a/libclc/opencl/include/clc/opencl/math/nan.inc b/libclc/opencl/include/clc/opencl/math/nan.inc
index 45b96f8..3134e56 100644
--- a/libclc/opencl/include/clc/opencl/math/nan.inc
+++ b/libclc/opencl/include/clc/opencl/math/nan.inc
@@ -6,4 +6,4 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE nan(__CLC_U_GENTYPE code);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE nan(__CLC_U_GENTYPE code);
diff --git a/libclc/opencl/include/clc/opencl/relational/all.h b/libclc/opencl/include/clc/opencl/relational/all.h
index 6c19d83..f7f1400 100644
--- a/libclc/opencl/include/clc/opencl/relational/all.h
+++ b/libclc/opencl/include/clc/opencl/relational/all.h
@@ -11,7 +11,7 @@
#include <clc/opencl/opencl-base.h>
-#define _CLC_ALL_DECL(TYPE) _CLC_OVERLOAD _CLC_DECL int all(TYPE v);
+#define _CLC_ALL_DECL(TYPE) _CLC_OVERLOAD _CLC_CONST _CLC_DECL int all(TYPE v);
#define _CLC_VECTOR_ALL_DECL(TYPE) \
_CLC_ALL_DECL(TYPE) \
diff --git a/libclc/opencl/include/clc/opencl/relational/any.h b/libclc/opencl/include/clc/opencl/relational/any.h
index 4f6b2e2..7b1ad33 100644
--- a/libclc/opencl/include/clc/opencl/relational/any.h
+++ b/libclc/opencl/include/clc/opencl/relational/any.h
@@ -11,7 +11,7 @@
#include <clc/opencl/opencl-base.h>
-#define _CLC_ANY_DECL(TYPE) _CLC_OVERLOAD _CLC_DECL int any(TYPE v);
+#define _CLC_ANY_DECL(TYPE) _CLC_OVERLOAD _CLC_CONST _CLC_DECL int any(TYPE v);
#define _CLC_VECTOR_ANY_DECL(TYPE) \
_CLC_ANY_DECL(TYPE) \
diff --git a/libclc/opencl/include/clc/opencl/relational/bitselect.inc b/libclc/opencl/include/clc/opencl/relational/bitselect.inc
index d821bb8..429ef37 100644
--- a/libclc/opencl/include/clc/opencl/relational/bitselect.inc
+++ b/libclc/opencl/include/clc/opencl/relational/bitselect.inc
@@ -6,6 +6,6 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE bitselect(__CLC_GENTYPE x,
- __CLC_GENTYPE y,
- __CLC_GENTYPE z);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE bitselect(__CLC_GENTYPE x,
+ __CLC_GENTYPE y,
+ __CLC_GENTYPE z);
diff --git a/libclc/opencl/include/clc/opencl/relational/isequal.h b/libclc/opencl/include/clc/opencl/relational/isequal.h
index c679b2c..b3dff2e 100644
--- a/libclc/opencl/include/clc/opencl/relational/isequal.h
+++ b/libclc/opencl/include/clc/opencl/relational/isequal.h
@@ -12,7 +12,7 @@
#include <clc/opencl/opencl-base.h>
#define _CLC_ISEQUAL_DECL(TYPE, RETTYPE) \
- _CLC_OVERLOAD _CLC_DECL RETTYPE isequal(TYPE x, TYPE y);
+ _CLC_OVERLOAD _CLC_CONST _CLC_DECL RETTYPE isequal(TYPE x, TYPE y);
#define _CLC_VECTOR_ISEQUAL_DECL(TYPE, RETTYPE) \
_CLC_ISEQUAL_DECL(TYPE##2, RETTYPE##2) \
diff --git a/libclc/opencl/include/clc/opencl/relational/isinf.h b/libclc/opencl/include/clc/opencl/relational/isinf.h
index 561cb1c..0dd7246 100644
--- a/libclc/opencl/include/clc/opencl/relational/isinf.h
+++ b/libclc/opencl/include/clc/opencl/relational/isinf.h
@@ -12,7 +12,7 @@
#include <clc/opencl/opencl-base.h>
#define _CLC_ISINF_DECL(RET_TYPE, ARG_TYPE) \
- _CLC_OVERLOAD _CLC_DECL RET_TYPE isinf(ARG_TYPE);
+ _CLC_OVERLOAD _CLC_CONST _CLC_DECL RET_TYPE isinf(ARG_TYPE);
#define _CLC_VECTOR_ISINF_DECL(RET_TYPE, ARG_TYPE) \
_CLC_ISINF_DECL(RET_TYPE##2, ARG_TYPE##2) \
diff --git a/libclc/opencl/include/clc/opencl/relational/isnan.h b/libclc/opencl/include/clc/opencl/relational/isnan.h
index fe1e572..f02ec5d 100644
--- a/libclc/opencl/include/clc/opencl/relational/isnan.h
+++ b/libclc/opencl/include/clc/opencl/relational/isnan.h
@@ -12,7 +12,7 @@
#include <clc/opencl/opencl-base.h>
#define _CLC_ISNAN_DECL(RET_TYPE, ARG_TYPE) \
- _CLC_OVERLOAD _CLC_DECL RET_TYPE isnan(ARG_TYPE);
+ _CLC_OVERLOAD _CLC_CONST _CLC_DECL RET_TYPE isnan(ARG_TYPE);
#define _CLC_VECTOR_ISNAN_DECL(RET_TYPE, ARG_TYPE) \
_CLC_ISNAN_DECL(RET_TYPE##2, ARG_TYPE##2) \
diff --git a/libclc/opencl/include/clc/opencl/shared/clamp.inc b/libclc/opencl/include/clc/opencl/shared/clamp.inc
index 75dec9c..74d2f70 100644
--- a/libclc/opencl/include/clc/opencl/shared/clamp.inc
+++ b/libclc/opencl/include/clc/opencl/shared/clamp.inc
@@ -6,11 +6,12 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE clamp(__CLC_GENTYPE x, __CLC_GENTYPE y,
- __CLC_GENTYPE z);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE clamp(__CLC_GENTYPE x,
+ __CLC_GENTYPE y,
+ __CLC_GENTYPE z);
#ifndef __CLC_SCALAR
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE clamp(__CLC_GENTYPE x,
- __CLC_SCALAR_GENTYPE y,
- __CLC_SCALAR_GENTYPE z);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE clamp(__CLC_GENTYPE x,
+ __CLC_SCALAR_GENTYPE y,
+ __CLC_SCALAR_GENTYPE z);
#endif
diff --git a/libclc/opencl/include/clc/opencl/shared/max.inc b/libclc/opencl/include/clc/opencl/shared/max.inc
index 98610df..99c28a7 100644
--- a/libclc/opencl/include/clc/opencl/shared/max.inc
+++ b/libclc/opencl/include/clc/opencl/shared/max.inc
@@ -6,9 +6,10 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE max(__CLC_GENTYPE a, __CLC_GENTYPE b);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE max(__CLC_GENTYPE a,
+ __CLC_GENTYPE b);
#ifndef __CLC_SCALAR
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE max(__CLC_GENTYPE a,
- __CLC_SCALAR_GENTYPE b);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE max(__CLC_GENTYPE a,
+ __CLC_SCALAR_GENTYPE b);
#endif
diff --git a/libclc/opencl/include/clc/opencl/shared/min.inc b/libclc/opencl/include/clc/opencl/shared/min.inc
index 3140b63..ce5d906 100644
--- a/libclc/opencl/include/clc/opencl/shared/min.inc
+++ b/libclc/opencl/include/clc/opencl/shared/min.inc
@@ -6,9 +6,10 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE min(__CLC_GENTYPE a, __CLC_GENTYPE b);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE min(__CLC_GENTYPE a,
+ __CLC_GENTYPE b);
#ifndef __CLC_SCALAR
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE min(__CLC_GENTYPE a,
- __CLC_SCALAR_GENTYPE b);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL __CLC_GENTYPE min(__CLC_GENTYPE a,
+ __CLC_SCALAR_GENTYPE b);
#endif
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_global_id.h b/libclc/opencl/include/clc/opencl/workitem/get_global_id.h
index 98217c0..d64da9d 100644
--- a/libclc/opencl/include/clc/opencl/workitem/get_global_id.h
+++ b/libclc/opencl/include/clc/opencl/workitem/get_global_id.h
@@ -11,6 +11,6 @@
#include <clc/opencl/opencl-base.h>
-_CLC_DECL _CLC_OVERLOAD size_t get_global_id(uint dim);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL size_t get_global_id(uint dim);
#endif // __CLC_OPENCL_WORKITEM_GET_GLOBAL_ID_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_global_offset.h b/libclc/opencl/include/clc/opencl/workitem/get_global_offset.h
index 8c82033..ac9704e 100644
--- a/libclc/opencl/include/clc/opencl/workitem/get_global_offset.h
+++ b/libclc/opencl/include/clc/opencl/workitem/get_global_offset.h
@@ -11,6 +11,6 @@
#include <clc/opencl/opencl-base.h>
-_CLC_DECL _CLC_OVERLOAD size_t get_global_offset(uint dim);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL size_t get_global_offset(uint dim);
#endif // __CLC_OPENCL_WORKITEM_GET_GLOBAL_OFFSET_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_global_size.h b/libclc/opencl/include/clc/opencl/workitem/get_global_size.h
index 5c74260..4809cc4b 100644
--- a/libclc/opencl/include/clc/opencl/workitem/get_global_size.h
+++ b/libclc/opencl/include/clc/opencl/workitem/get_global_size.h
@@ -11,6 +11,6 @@
#include <clc/opencl/opencl-base.h>
-_CLC_DECL _CLC_OVERLOAD size_t get_global_size(uint dim);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL size_t get_global_size(uint dim);
#endif // __CLC_OPENCL_WORKITEM_GET_GLOBAL_SIZE_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_group_id.h b/libclc/opencl/include/clc/opencl/workitem/get_group_id.h
index 3d0f51a..e3d8b07 100644
--- a/libclc/opencl/include/clc/opencl/workitem/get_group_id.h
+++ b/libclc/opencl/include/clc/opencl/workitem/get_group_id.h
@@ -11,6 +11,6 @@
#include <clc/opencl/opencl-base.h>
-_CLC_DECL _CLC_OVERLOAD size_t get_group_id(uint dim);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL size_t get_group_id(uint dim);
#endif // __CLC_OPENCL_WORKITEM_GET_GROUP_ID_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_local_id.h b/libclc/opencl/include/clc/opencl/workitem/get_local_id.h
index 31fe9d2..88fbad98 100644
--- a/libclc/opencl/include/clc/opencl/workitem/get_local_id.h
+++ b/libclc/opencl/include/clc/opencl/workitem/get_local_id.h
@@ -11,6 +11,6 @@
#include <clc/opencl/opencl-base.h>
-_CLC_DECL _CLC_OVERLOAD size_t get_local_id(uint dim);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL size_t get_local_id(uint dim);
#endif // __CLC_OPENCL_WORKITEM_GET_LOCAL_ID_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_local_linear_id.h b/libclc/opencl/include/clc/opencl/workitem/get_local_linear_id.h
index 9a32aa4..072318d 100644
--- a/libclc/opencl/include/clc/opencl/workitem/get_local_linear_id.h
+++ b/libclc/opencl/include/clc/opencl/workitem/get_local_linear_id.h
@@ -11,6 +11,6 @@
#include <clc/opencl/opencl-base.h>
-_CLC_OVERLOAD _CLC_DECL size_t get_local_linear_id();
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL size_t get_local_linear_id();
#endif // __CLC_OPENCL_WORKITEM_GET_LOCAL_LINEAR_ID_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_local_size.h b/libclc/opencl/include/clc/opencl/workitem/get_local_size.h
index c93a7a2..49077af 100644
--- a/libclc/opencl/include/clc/opencl/workitem/get_local_size.h
+++ b/libclc/opencl/include/clc/opencl/workitem/get_local_size.h
@@ -11,6 +11,6 @@
#include <clc/opencl/opencl-base.h>
-_CLC_DECL _CLC_OVERLOAD size_t get_local_size(uint dim);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL size_t get_local_size(uint dim);
#endif // __CLC_OPENCL_WORKITEM_GET_LOCAL_SIZE_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_max_sub_group_size.h b/libclc/opencl/include/clc/opencl/workitem/get_max_sub_group_size.h
index 140f224..b127376 100644
--- a/libclc/opencl/include/clc/opencl/workitem/get_max_sub_group_size.h
+++ b/libclc/opencl/include/clc/opencl/workitem/get_max_sub_group_size.h
@@ -11,6 +11,6 @@
#include <clc/opencl/opencl-base.h>
-_CLC_OVERLOAD _CLC_DECL uint get_max_sub_group_size();
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL uint get_max_sub_group_size();
#endif // __CLC_OPENCL_WORKITEM_GET_MAX_SUB_GROUP_SIZE_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_num_groups.h b/libclc/opencl/include/clc/opencl/workitem/get_num_groups.h
index 0c47a5e..2978316 100644
--- a/libclc/opencl/include/clc/opencl/workitem/get_num_groups.h
+++ b/libclc/opencl/include/clc/opencl/workitem/get_num_groups.h
@@ -11,6 +11,6 @@
#include <clc/opencl/opencl-base.h>
-_CLC_DECL _CLC_OVERLOAD size_t get_num_groups(uint dim);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL size_t get_num_groups(uint dim);
#endif // __CLC_OPENCL_WORKITEM_GET_NUM_GROUPS_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_num_sub_groups.h b/libclc/opencl/include/clc/opencl/workitem/get_num_sub_groups.h
index 5c72db0..f9fd2cd 100644
--- a/libclc/opencl/include/clc/opencl/workitem/get_num_sub_groups.h
+++ b/libclc/opencl/include/clc/opencl/workitem/get_num_sub_groups.h
@@ -11,6 +11,6 @@
#include <clc/opencl/opencl-base.h>
-_CLC_OVERLOAD _CLC_DECL uint get_num_sub_groups();
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL uint get_num_sub_groups();
#endif // __CLC_OPENCL_WORKITEM_GET_NUM_SUB_GROUPS_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_sub_group_id.h b/libclc/opencl/include/clc/opencl/workitem/get_sub_group_id.h
index 3dbf39b..0a67c02 100644
--- a/libclc/opencl/include/clc/opencl/workitem/get_sub_group_id.h
+++ b/libclc/opencl/include/clc/opencl/workitem/get_sub_group_id.h
@@ -11,6 +11,6 @@
#include <clc/opencl/opencl-base.h>
-_CLC_OVERLOAD _CLC_DECL uint get_sub_group_id();
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL uint get_sub_group_id();
#endif // __CLC_OPENCL_WORKITEM_GET_SUB_GROUP_ID_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_sub_group_local_id.h b/libclc/opencl/include/clc/opencl/workitem/get_sub_group_local_id.h
index 7e139a8..37891d8 100644
--- a/libclc/opencl/include/clc/opencl/workitem/get_sub_group_local_id.h
+++ b/libclc/opencl/include/clc/opencl/workitem/get_sub_group_local_id.h
@@ -11,6 +11,6 @@
#include <clc/opencl/opencl-base.h>
-_CLC_OVERLOAD _CLC_DECL uint get_sub_group_local_id();
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL uint get_sub_group_local_id();
#endif // __CLC_OPENCL_WORKITEM_GET_SUB_GROUP_LOCAL_ID_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_sub_group_size.h b/libclc/opencl/include/clc/opencl/workitem/get_sub_group_size.h
index dd4f89d..74e9b03 100644
--- a/libclc/opencl/include/clc/opencl/workitem/get_sub_group_size.h
+++ b/libclc/opencl/include/clc/opencl/workitem/get_sub_group_size.h
@@ -11,6 +11,6 @@
#include <clc/opencl/opencl-base.h>
-_CLC_OVERLOAD _CLC_DECL uint get_sub_group_size();
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL uint get_sub_group_size();
#endif // __CLC_OPENCL_WORKITEM_GET_SUB_GROUP_SIZE_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_work_dim.h b/libclc/opencl/include/clc/opencl/workitem/get_work_dim.h
index c4c3c1b..344ffaa 100644
--- a/libclc/opencl/include/clc/opencl/workitem/get_work_dim.h
+++ b/libclc/opencl/include/clc/opencl/workitem/get_work_dim.h
@@ -11,6 +11,6 @@
#include <clc/opencl/opencl-base.h>
-_CLC_DECL _CLC_OVERLOAD uint get_work_dim(void);
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL uint get_work_dim(void);
#endif // __CLC_OPENCL_WORKITEM_GET_WORK_DIM_H__