[libclc] [libclc] Add __attribute__((const)) to functions that don't access memory (PR #152456)
Wenju He via cfe-commits
cfe-commits at lists.llvm.org
Thu Aug 7 01:35:50 PDT 2025
https://github.com/wenju-he created https://github.com/llvm/llvm-project/pull/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.
>From 6cc621cdda4b873ab1cebc92707553a760630955 Mon Sep 17 00:00:00 2001
From: Wenju He <wenju.he at intel.com>
Date: Thu, 7 Aug 2025 09:27:26 +0200
Subject: [PATCH] [libclc] Add __attribute__((const)) to functions that don't
access memory
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.
---
libclc/clc/include/clc/clcfunc.h | 1 +
libclc/clc/include/clc/common/clc_smoothstep.inc | 1 +
libclc/clc/include/clc/geometric/binary_decl.inc | 4 ++--
libclc/clc/include/clc/geometric/clc_cross.h | 12 ++++++------
libclc/clc/include/clc/geometric/unary_decl.inc | 2 +-
libclc/clc/include/clc/integer/clc_abs.inc | 2 +-
libclc/clc/include/clc/integer/clc_abs_diff.inc | 4 ++--
.../clc/integer/clc_bitfield_extract_decl.inc | 4 ++--
.../include/clc/integer/clc_bitfield_insert.inc | 7 ++++---
.../math/binary_decl_with_scalar_second_arg.inc | 8 ++++----
.../clc/include/clc/math/binary_def_via_fp32.inc | 4 ++--
libclc/clc/include/clc/math/clc_exp_helper.inc | 8 +++-----
libclc/clc/include/clc/math/clc_ldexp.inc | 3 ++-
libclc/clc/include/clc/math/clc_nan.inc | 4 ++--
libclc/clc/include/clc/math/unary_decl.inc | 3 +++
.../clc/math/unary_decl_with_int_return.inc | 2 +-
libclc/clc/include/clc/misc/shuffle2_decl.inc | 8 ++++----
libclc/clc/include/clc/misc/shuffle_decl.inc | 8 ++++----
.../clc/include/clc/relational/binary_decl.inc | 3 ++-
libclc/clc/include/clc/relational/clc_all.h | 3 ++-
libclc/clc/include/clc/relational/clc_any.h | 3 ++-
.../clc/include/clc/relational/clc_bitselect.inc | 5 ++---
libclc/clc/include/clc/relational/clc_isequal.h | 2 +-
libclc/clc/include/clc/relational/clc_isinf.h | 2 +-
libclc/clc/include/clc/relational/clc_isnan.h | 2 +-
.../include/clc/relational/clc_select_decl.inc | 10 ++++------
libclc/clc/include/clc/relational/unary_decl.inc | 2 +-
libclc/clc/include/clc/shared/binary_decl.inc | 3 +++
.../shared/binary_decl_with_int_second_arg.inc | 3 +++
libclc/clc/include/clc/shared/clc_clamp.inc | 11 +++++------
libclc/clc/include/clc/shared/clc_max.inc | 8 ++++----
libclc/clc/include/clc/shared/clc_min.inc | 8 ++++----
libclc/clc/include/clc/shared/ternary_decl.inc | 3 +++
libclc/clc/include/clc/shared/unary_decl.inc | 3 +++
.../clc/include/clc/workitem/clc_get_global_id.h | 2 +-
.../include/clc/workitem/clc_get_global_offset.h | 2 +-
.../include/clc/workitem/clc_get_global_size.h | 2 +-
.../clc/include/clc/workitem/clc_get_group_id.h | 2 +-
.../clc/include/clc/workitem/clc_get_local_id.h | 2 +-
.../clc/workitem/clc_get_local_linear_id.h | 2 +-
.../include/clc/workitem/clc_get_local_size.h | 2 +-
.../clc/workitem/clc_get_max_sub_group_size.h | 2 +-
.../include/clc/workitem/clc_get_num_groups.h | 2 +-
.../clc/workitem/clc_get_num_sub_groups.h | 2 +-
.../include/clc/workitem/clc_get_sub_group_id.h | 2 +-
.../clc/workitem/clc_get_sub_group_local_id.h | 2 +-
.../clc/workitem/clc_get_sub_group_size.h | 2 +-
.../clc/include/clc/workitem/clc_get_work_dim.h | 2 +-
libclc/opencl/include/clc/opencl/common/mix.inc | 10 ++++++----
.../include/clc/opencl/common/smoothstep.inc | 16 +++++++++-------
libclc/opencl/include/clc/opencl/common/step.inc | 7 ++++---
.../opencl/include/clc/opencl/geometric/cross.h | 8 ++++----
libclc/opencl/include/clc/opencl/integer/abs.inc | 2 +-
.../include/clc/opencl/integer/abs_diff.inc | 4 ++--
.../opencl/include/clc/opencl/integer/upsample.h | 2 +-
libclc/opencl/include/clc/opencl/math/ldexp.inc | 2 +-
libclc/opencl/include/clc/opencl/math/nan.inc | 2 +-
.../opencl/include/clc/opencl/relational/all.h | 2 +-
.../opencl/include/clc/opencl/relational/any.h | 2 +-
.../include/clc/opencl/relational/bitselect.inc | 6 +++---
.../include/clc/opencl/relational/isequal.h | 2 +-
.../opencl/include/clc/opencl/relational/isinf.h | 2 +-
.../opencl/include/clc/opencl/relational/isnan.h | 2 +-
.../opencl/include/clc/opencl/shared/clamp.inc | 11 ++++++-----
libclc/opencl/include/clc/opencl/shared/max.inc | 7 ++++---
libclc/opencl/include/clc/opencl/shared/min.inc | 7 ++++---
.../include/clc/opencl/workitem/get_global_id.h | 2 +-
.../clc/opencl/workitem/get_global_offset.h | 2 +-
.../clc/opencl/workitem/get_global_size.h | 2 +-
.../include/clc/opencl/workitem/get_group_id.h | 2 +-
.../include/clc/opencl/workitem/get_local_id.h | 2 +-
.../clc/opencl/workitem/get_local_linear_id.h | 2 +-
.../include/clc/opencl/workitem/get_local_size.h | 2 +-
.../clc/opencl/workitem/get_max_sub_group_size.h | 2 +-
.../include/clc/opencl/workitem/get_num_groups.h | 2 +-
.../clc/opencl/workitem/get_num_sub_groups.h | 2 +-
.../clc/opencl/workitem/get_sub_group_id.h | 2 +-
.../clc/opencl/workitem/get_sub_group_local_id.h | 2 +-
.../clc/opencl/workitem/get_sub_group_size.h | 2 +-
.../include/clc/opencl/workitem/get_work_dim.h | 2 +-
80 files changed, 165 insertions(+), 141 deletions(-)
diff --git a/libclc/clc/include/clc/clcfunc.h b/libclc/clc/include/clc/clcfunc.h
index 10b9cdd099900..b09748f3bf0d9 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_CNFN __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 158c578d73d95..9406570fff0b1 100644
--- a/libclc/clc/include/clc/common/clc_smoothstep.inc
+++ b/libclc/clc/include/clc/common/clc_smoothstep.inc
@@ -6,6 +6,7 @@
//
//===----------------------------------------------------------------------===//
+_CLC_CNFN
_CLC_OVERLOAD _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 342d97a802659..2118c070ed867 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_CNFN _CLC_OVERLOAD _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 e5aa913abfa29..701fe88502699 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_CNFN _CLC_OVERLOAD _CLC_DECL float3 __clc_cross(float3 p0, float3 p1);
+_CLC_CNFN _CLC_OVERLOAD _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_CNFN _CLC_OVERLOAD _CLC_DECL double3 __clc_cross(double3 p0, double3 p1);
+_CLC_CNFN _CLC_OVERLOAD _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_CNFN _CLC_OVERLOAD _CLC_DECL half3 __clc_cross(half3 p0, half3 p1);
+_CLC_CNFN _CLC_OVERLOAD _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 cff66000fcf77..0de1a0924bdcd 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_CNFN _CLC_OVERLOAD _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 f1e2c5a5149cb..dc84a605485b1 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_CNFN _CLC_OVERLOAD _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 2514a37b78144..9cc9f7aea380f 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_CNFN _CLC_OVERLOAD _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 c93eff08de0bc..ea971480fd958 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_CNFN _CLC_OVERLOAD _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 22f58bdc09830..0981f63b34f71 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_CNFN _CLC_OVERLOAD _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 1b2b9ff33521b..19d821587df37 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_CNFN _CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_GENTYPE a,
+ __CLC_GENTYPE b);
+_CLC_CNFN _CLC_OVERLOAD _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 b53a178bf0f10..84bc85265d48b 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_CNFN _CLC_OVERLOAD _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 cdf650405c815..2ba7f7143cf2b 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_CNFN _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);
diff --git a/libclc/clc/include/clc/math/clc_ldexp.inc b/libclc/clc/include/clc/math/clc_ldexp.inc
index 9b88359d212d2..cb86f30369586 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_CNFN _CLC_DECL _CLC_OVERLOAD __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 22604e0a5785e..a26e959fd4dc1 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_CNFN _CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __clc_nan(__CLC_U_GENTYPE code);
+_CLC_CNFN _CLC_OVERLOAD _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 ce1a71b8284aa..6d822cca61276 100644
--- a/libclc/clc/include/clc/math/unary_decl.inc
+++ b/libclc/clc/include/clc/math/unary_decl.inc
@@ -6,4 +6,7 @@
//
//===----------------------------------------------------------------------===//
+#ifndef _CLC_NOT_CNFN
+_CLC_CNFN
+#endif
_CLC_OVERLOAD _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 ed8d0a12506f2..9e36c086a2dc3 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_CNFN _CLC_OVERLOAD _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 47fd5e28a3142..9787cce320a2d 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_CNFN _CLC_OVERLOAD _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_CNFN _CLC_OVERLOAD _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_CNFN _CLC_OVERLOAD _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_CNFN _CLC_OVERLOAD _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 c6c64d5b5fede..cd3c51bf7db5f 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_CNFN _CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 2) x, __CLC_U_GENTYPE mask);
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
+_CLC_CNFN _CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 4) x, __CLC_U_GENTYPE mask);
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
+_CLC_CNFN _CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 8) x, __CLC_U_GENTYPE mask);
-_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
+_CLC_CNFN _CLC_OVERLOAD _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 dc8ec9db7b7da..c4968c207f153 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_CNFN _CLC_OVERLOAD _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 4d01fb7795b72..84bccdef340f0 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_CNFN _CLC_OVERLOAD _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 1e287af58d880..7536ed8ab8dba 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_CNFN _CLC_OVERLOAD _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 b81cfc2ef0468..a902d6ea30837 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_CNFN _CLC_OVERLOAD _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 f3789f595637e..b71e4b1730a50 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_CNFN _CLC_OVERLOAD _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 39a9de6c434e1..aa357e98c87e3 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_CNFN _CLC_OVERLOAD _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 d0821a3c52639..a6bbd50d2f7fd 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_CNFN _CLC_OVERLOAD _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 1cf026dc8db36..c8d6f41588f10 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_CNFN _CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
+__CLC_SELECT_FN(__CLC_GENTYPE x, __CLC_GENTYPE y, __CLC_S_GENTYPE z);
+_CLC_CNFN _CLC_OVERLOAD _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 cc3f2d065529b..961497ad4278d 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_CNFN _CLC_OVERLOAD _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 ff4739dcdb8d1..7989e7882e7e0 100644
--- a/libclc/clc/include/clc/shared/binary_decl.inc
+++ b/libclc/clc/include/clc/shared/binary_decl.inc
@@ -6,5 +6,8 @@
//
//===----------------------------------------------------------------------===//
+#ifndef _CLC_NOT_CNFN
+_CLC_CNFN
+#endif
_CLC_OVERLOAD _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 a86d89b4fc85f..d8fafd904278c 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,7 @@
#include <clc/utils.h>
+#ifndef _CLC_NOT_CNFN
+_CLC_CNFN
+#endif
_CLC_OVERLOAD _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 c1280a6f699ac..e530d23051e54 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_CNFN _CLC_OVERLOAD _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_CNFN _CLC_OVERLOAD _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 ca99c1264ad62..dab1ffe2811e1 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_CNFN _CLC_OVERLOAD _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_CNFN _CLC_OVERLOAD _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 c7cfe23ec3cd6..c3670de99a5f2 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_CNFN _CLC_OVERLOAD _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_CNFN _CLC_OVERLOAD _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 a76db1eefb6a4..57200ca3bb997 100644
--- a/libclc/clc/include/clc/shared/ternary_decl.inc
+++ b/libclc/clc/include/clc/shared/ternary_decl.inc
@@ -6,5 +6,8 @@
//
//===----------------------------------------------------------------------===//
+#ifndef _CLC_NOT_CNFN
+_CLC_CNFN
+#endif
_CLC_OVERLOAD _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 ce1a71b8284aa..6d822cca61276 100644
--- a/libclc/clc/include/clc/shared/unary_decl.inc
+++ b/libclc/clc/include/clc/shared/unary_decl.inc
@@ -6,4 +6,7 @@
//
//===----------------------------------------------------------------------===//
+#ifndef _CLC_NOT_CNFN
+_CLC_CNFN
+#endif
_CLC_OVERLOAD _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 da15ed8e53ca8..eed823dd3f691 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_CNFN _CLC_OVERLOAD _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 639be48d0b958..41d837701638a 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_CNFN _CLC_OVERLOAD _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 fce11f417bf9d..35ffb2baa5a42 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_CNFN _CLC_OVERLOAD _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 3fd43cb77f893..35597f606c918 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_CNFN _CLC_OVERLOAD _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 c446127a5f90c..eda16794f1e79 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_CNFN _CLC_OVERLOAD _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 34c47f2afcb30..09cc54f960927 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_CNFN _CLC_OVERLOAD _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 7fe3c193e0ec1..91ee7220b8c4a 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_CNFN _CLC_OVERLOAD _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 8b521fe32862d..9ec5b555c4ba0 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_CNFN _CLC_OVERLOAD _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 f860944ad517f..bb4a189df9aaa 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_CNFN _CLC_OVERLOAD _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 6965aef1bce30..5bf64cfd58101 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_CNFN _CLC_OVERLOAD _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 ac3d7bd30e454..7e97ab242d3ea 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_CNFN _CLC_OVERLOAD _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 06bb6f8b77a5c..ec3085a04234b 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_CNFN _CLC_OVERLOAD _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 a6e8e49470e4f..330b1603f4cd7 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_CNFN _CLC_OVERLOAD _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 7e2da4570c4f5..dd23dd014ce60 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_CNFN _CLC_OVERLOAD _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 2b8350f5baae1..ab5f089d3403c 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_CNFN _CLC_OVERLOAD _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_CNFN _CLC_OVERLOAD _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 4547483364ba2..5046049a0611e 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_CNFN _CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE smoothstep(__CLC_GENTYPE edge0,
+ __CLC_GENTYPE edge1,
+ __CLC_GENTYPE x);
+_CLC_CNFN _CLC_OVERLOAD _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_CNFN _CLC_OVERLOAD _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 86d4ed03dbcab..6145e0e70d934 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_CNFN _CLC_OVERLOAD _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_CNFN _CLC_OVERLOAD _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 be137965aea4b..eca140c5ea0b6 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_CNFN _CLC_OVERLOAD _CLC_DECL float3 cross(float3 p0, float3 p1);
+_CLC_CNFN _CLC_OVERLOAD _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_CNFN _CLC_OVERLOAD _CLC_DECL double3 cross(double3 p0, double3 p1);
+_CLC_CNFN _CLC_OVERLOAD _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 6620ce36f8c6c..f4ead301d1321 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_CNFN _CLC_OVERLOAD _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 41873854021eb..4ca0ddf98bba1 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_CNFN _CLC_OVERLOAD _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 a713f2b392f27..6d13f350605f3 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_CNFN _CLC_OVERLOAD _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 b5a5cfcafdd53..748289c90ec8f 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_CNFN _CLC_OVERLOAD _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 45b96f81d7285..73817ea9edc0f 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_CNFN _CLC_OVERLOAD _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 6c19d83e25bbf..ccce892c41c98 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_CNFN _CLC_OVERLOAD _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 4f6b2e22460c1..bd7d1ae210976 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_CNFN _CLC_OVERLOAD _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 d821bb86a1409..79191606b0027 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_CNFN _CLC_OVERLOAD _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 c679b2cfc3feb..7fc37d4cef0f5 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_CNFN _CLC_OVERLOAD _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 561cb1cda47b7..eb782e2fa6e22 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_CNFN _CLC_OVERLOAD _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 fe1e5721c1d24..13130942a6d80 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_CNFN _CLC_OVERLOAD _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 75dec9c258e8a..3b99a4e5beb89 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_CNFN _CLC_OVERLOAD _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_CNFN _CLC_OVERLOAD _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 98610dfdb1018..718c17552bcac 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_CNFN _CLC_OVERLOAD _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_CNFN _CLC_OVERLOAD _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 3140b637f3953..abbcd6c61cc0a 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_CNFN _CLC_OVERLOAD _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_CNFN _CLC_OVERLOAD _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 98217c0ba6b4d..1993f47fc1e77 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_CNFN _CLC_OVERLOAD _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 8c820331e6ab2..1046b31b9f5e2 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_CNFN _CLC_OVERLOAD _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 5c7426047d7d0..bdaafa98ecd9c 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_CNFN _CLC_OVERLOAD _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 3d0f51aeb8dc2..88350fc6d9575 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_CNFN _CLC_OVERLOAD _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 31fe9d2943f48..625add8a2f0f8 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_CNFN _CLC_OVERLOAD _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 9a32aa46b8191..5bcfbc7f988cb 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_CNFN _CLC_OVERLOAD _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 c93a7a27bd7cf..f3fb360a3b718 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_CNFN _CLC_OVERLOAD _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 140f22426ef6b..f9a8c4d0e764c 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_CNFN _CLC_OVERLOAD _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 0c47a5ef462b3..cdd64b47a90bd 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_CNFN _CLC_OVERLOAD _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 5c72db0b40c72..5cb995fa086c0 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_CNFN _CLC_OVERLOAD _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 3dbf39b8cfbad..cab7bbfae530f 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_CNFN _CLC_OVERLOAD _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 7e139a857a7d7..1c6283f21f75f 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_CNFN _CLC_OVERLOAD _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 dd4f89da18a14..566ac6a83cb2d 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_CNFN _CLC_OVERLOAD _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 c4c3c1b224a9a..c298c65e09176 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_CNFN _CLC_OVERLOAD _CLC_DECL uint get_work_dim(void);
#endif // __CLC_OPENCL_WORKITEM_GET_WORK_DIM_H__
More information about the cfe-commits
mailing list