haonanya created this revision.
haonanya added reviewers: Anastasia, svenvh.
haonanya added a project: clang.
Herald added subscribers: Naghasan, ldrumm, yaxunl.
haonanya requested review of this revision.
Herald added a subscriber: cfe-commits.
It is necessary to guard atomic_double type according to
https://www.khronos.org/registry/OpenCL/specs/3.0-unified/html/OpenCL_C.html#_footnotedef_54.
Platform that disables cl_khr_int64_base_atomics and
cl_khr_int64_extended_atomics will have compiling errors even if atomic_double
is not used.
Compile the following OpenCL test with options clang++ -cc1 -emit-llvm -triple
spir-unknown-unknown -finclude-default-header
-cl-ext=+cl_khr_fp64,+__opencl_c_fp64,-cl_khr_int64_base_atomics -cl-std=CL3.0
test.cl
__kernel void test_kernel1() {}
use of type 'atomic_double' (aka '_Atomic(double)') requires
cl_khr_int64_extended_atomics extension to be enabled
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D119398
Files:
clang/lib/Headers/opencl-c.h
clang/lib/Sema/OpenCLBuiltins.td
Index: clang/lib/Sema/OpenCLBuiltins.td
===================================================================
--- clang/lib/Sema/OpenCLBuiltins.td
+++ clang/lib/Sema/OpenCLBuiltins.td
@@ -95,22 +95,22 @@
def FuncExtFloatAtomicsFp16GenericLoadStore :
FunctionExtension<"cl_ext_float_atomics
__opencl_c_ext_fp16_global_atomic_load_store
__opencl_c_ext_fp16_local_atomic_load_store">;
def FuncExtFloatAtomicsFp16GlobalAdd :
FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_global_atomic_add">;
def FuncExtFloatAtomicsFp32GlobalAdd :
FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_global_atomic_add">;
-def FuncExtFloatAtomicsFp64GlobalAdd :
FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_global_atomic_add">;
+def FuncExtFloatAtomicsFp64GlobalAdd :
FunctionExtension<"cl_ext_float_atomics cl_khr_int64_base_atomics
cl_khr_int64_extended_atomics __opencl_c_ext_fp64_global_atomic_add">;
def FuncExtFloatAtomicsFp16LocalAdd :
FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_local_atomic_add">;
def FuncExtFloatAtomicsFp32LocalAdd :
FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_local_atomic_add">;
-def FuncExtFloatAtomicsFp64LocalAdd :
FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_local_atomic_add">;
+def FuncExtFloatAtomicsFp64LocalAdd :
FunctionExtension<"cl_ext_float_atomics cl_khr_int64_base_atomics
cl_khr_int64_extended_atomics __opencl_c_ext_fp64_local_atomic_add">;
def FuncExtFloatAtomicsFp16GenericAdd :
FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_local_atomic_add
__opencl_c_ext_fp16_global_atomic_add">;
def FuncExtFloatAtomicsFp32GenericAdd :
FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_local_atomic_add
__opencl_c_ext_fp32_global_atomic_add">;
-def FuncExtFloatAtomicsFp64GenericAdd :
FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_local_atomic_add
__opencl_c_ext_fp64_global_atomic_add">;
+def FuncExtFloatAtomicsFp64GenericAdd :
FunctionExtension<"cl_ext_float_atomics cl_khr_int64_base_atomics
cl_khr_int64_extended_atomics __opencl_c_ext_fp64_local_atomic_add
__opencl_c_ext_fp64_global_atomic_add">;
def FuncExtFloatAtomicsFp16GlobalMinMax :
FunctionExtension<"cl_ext_float_atomics
__opencl_c_ext_fp16_global_atomic_min_max">;
def FuncExtFloatAtomicsFp32GlobalMinMax :
FunctionExtension<"cl_ext_float_atomics
__opencl_c_ext_fp32_global_atomic_min_max">;
-def FuncExtFloatAtomicsFp64GlobalMinMax :
FunctionExtension<"cl_ext_float_atomics
__opencl_c_ext_fp64_global_atomic_min_max">;
+def FuncExtFloatAtomicsFp64GlobalMinMax :
FunctionExtension<"cl_ext_float_atomics cl_khr_int64_base_atomics
cl_khr_int64_extended_atomics __opencl_c_ext_fp64_global_atomic_min_max">;
def FuncExtFloatAtomicsFp16LocalMinMax :
FunctionExtension<"cl_ext_float_atomics
__opencl_c_ext_fp16_local_atomic_min_max">;
def FuncExtFloatAtomicsFp32LocalMinMax :
FunctionExtension<"cl_ext_float_atomics
__opencl_c_ext_fp32_local_atomic_min_max">;
-def FuncExtFloatAtomicsFp64LocalMinMax :
FunctionExtension<"cl_ext_float_atomics
__opencl_c_ext_fp64_local_atomic_min_max">;
+def FuncExtFloatAtomicsFp64LocalMinMax :
FunctionExtension<"cl_ext_float_atomics cl_khr_int64_base_atomics
cl_khr_int64_extended_atomics __opencl_c_ext_fp64_local_atomic_min_max">;
def FuncExtFloatAtomicsFp16GenericMinMax :
FunctionExtension<"cl_ext_float_atomics
__opencl_c_ext_fp16_local_atomic_min_max
__opencl_c_ext_fp16_global_atomic_min_max">;
def FuncExtFloatAtomicsFp32GenericMinMax :
FunctionExtension<"cl_ext_float_atomics
__opencl_c_ext_fp32_local_atomic_min_max
__opencl_c_ext_fp32_global_atomic_min_max">;
-def FuncExtFloatAtomicsFp64GenericMinMax :
FunctionExtension<"cl_ext_float_atomics
__opencl_c_ext_fp64_local_atomic_min_max
__opencl_c_ext_fp64_global_atomic_min_max">;
+def FuncExtFloatAtomicsFp64GenericMinMax :
FunctionExtension<"cl_ext_float_atomics cl_khr_int64_base_atomics
cl_khr_int64_extended_atomics __opencl_c_ext_fp64_local_atomic_min_max
__opencl_c_ext_fp64_global_atomic_min_max">;
// Not a real extension, but a workaround to add C++ for OpenCL specific
builtins.
def FuncExtOpenCLCxx : FunctionExtension<"__cplusplus">;
Index: clang/lib/Headers/opencl-c.h
===================================================================
--- clang/lib/Headers/opencl-c.h
+++ clang/lib/Headers/opencl-c.h
@@ -13832,6 +13832,7 @@
#endif // defined(__opencl_c_ext_fp32_global_atomic_min_max) &&
\
defined(__opencl_c_ext_fp32_local_atomic_min_max)
+#if defined(cl_khr_int64_base_atomics) &&
defined(cl_khr_int64_extended_atomics)
#if defined(__opencl_c_ext_fp64_global_atomic_min_max)
double __ovld atomic_fetch_min(volatile __global atomic_double *object,
double operand);
@@ -13882,6 +13883,8 @@
memory_scope scope);
#endif // defined(__opencl_c_ext_fp64_global_atomic_min_max) &&
\
defined(__opencl_c_ext_fp64_local_atomic_min_max)
+#endif // defined(cl_khr_int64_base_atomics) &&
+ // defined(cl_khr_int64_extended_atomics)
#if defined(__opencl_c_ext_fp16_global_atomic_add)
half __ovld atomic_fetch_add(volatile __global atomic_half *object,
@@ -13985,6 +13988,7 @@
#endif // defined(__opencl_c_ext_fp32_global_atomic_add) &&
\
defined(__opencl_c_ext_fp32_local_atomic_add)
+#if defined(cl_khr_int64_base_atomics) &&
defined(cl_khr_int64_extended_atomics)
#if defined(__opencl_c_ext_fp64_global_atomic_add)
double __ovld atomic_fetch_add(volatile __global atomic_double *object,
double operand);
@@ -14035,6 +14039,8 @@
memory_scope scope);
#endif // defined(__opencl_c_ext_fp64_global_atomic_add) &&
\
defined(__opencl_c_ext_fp64_local_atomic_add)
+#endif // defined(cl_khr_int64_base_atomics) &&
+ // defined(cl_khr_int64_extended_atomics)
#endif // cl_ext_float_atomics
Index: clang/lib/Sema/OpenCLBuiltins.td
===================================================================
--- clang/lib/Sema/OpenCLBuiltins.td
+++ clang/lib/Sema/OpenCLBuiltins.td
@@ -95,22 +95,22 @@
def FuncExtFloatAtomicsFp16GenericLoadStore : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_global_atomic_load_store __opencl_c_ext_fp16_local_atomic_load_store">;
def FuncExtFloatAtomicsFp16GlobalAdd : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_global_atomic_add">;
def FuncExtFloatAtomicsFp32GlobalAdd : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_global_atomic_add">;
-def FuncExtFloatAtomicsFp64GlobalAdd : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_global_atomic_add">;
+def FuncExtFloatAtomicsFp64GlobalAdd : FunctionExtension<"cl_ext_float_atomics cl_khr_int64_base_atomics cl_khr_int64_extended_atomics __opencl_c_ext_fp64_global_atomic_add">;
def FuncExtFloatAtomicsFp16LocalAdd : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_local_atomic_add">;
def FuncExtFloatAtomicsFp32LocalAdd : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_local_atomic_add">;
-def FuncExtFloatAtomicsFp64LocalAdd : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_local_atomic_add">;
+def FuncExtFloatAtomicsFp64LocalAdd : FunctionExtension<"cl_ext_float_atomics cl_khr_int64_base_atomics cl_khr_int64_extended_atomics __opencl_c_ext_fp64_local_atomic_add">;
def FuncExtFloatAtomicsFp16GenericAdd : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_local_atomic_add __opencl_c_ext_fp16_global_atomic_add">;
def FuncExtFloatAtomicsFp32GenericAdd : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_local_atomic_add __opencl_c_ext_fp32_global_atomic_add">;
-def FuncExtFloatAtomicsFp64GenericAdd : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_local_atomic_add __opencl_c_ext_fp64_global_atomic_add">;
+def FuncExtFloatAtomicsFp64GenericAdd : FunctionExtension<"cl_ext_float_atomics cl_khr_int64_base_atomics cl_khr_int64_extended_atomics __opencl_c_ext_fp64_local_atomic_add __opencl_c_ext_fp64_global_atomic_add">;
def FuncExtFloatAtomicsFp16GlobalMinMax : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_global_atomic_min_max">;
def FuncExtFloatAtomicsFp32GlobalMinMax : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_global_atomic_min_max">;
-def FuncExtFloatAtomicsFp64GlobalMinMax : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_global_atomic_min_max">;
+def FuncExtFloatAtomicsFp64GlobalMinMax : FunctionExtension<"cl_ext_float_atomics cl_khr_int64_base_atomics cl_khr_int64_extended_atomics __opencl_c_ext_fp64_global_atomic_min_max">;
def FuncExtFloatAtomicsFp16LocalMinMax : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_local_atomic_min_max">;
def FuncExtFloatAtomicsFp32LocalMinMax : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_local_atomic_min_max">;
-def FuncExtFloatAtomicsFp64LocalMinMax : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_local_atomic_min_max">;
+def FuncExtFloatAtomicsFp64LocalMinMax : FunctionExtension<"cl_ext_float_atomics cl_khr_int64_base_atomics cl_khr_int64_extended_atomics __opencl_c_ext_fp64_local_atomic_min_max">;
def FuncExtFloatAtomicsFp16GenericMinMax : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_local_atomic_min_max __opencl_c_ext_fp16_global_atomic_min_max">;
def FuncExtFloatAtomicsFp32GenericMinMax : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_local_atomic_min_max __opencl_c_ext_fp32_global_atomic_min_max">;
-def FuncExtFloatAtomicsFp64GenericMinMax : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_local_atomic_min_max __opencl_c_ext_fp64_global_atomic_min_max">;
+def FuncExtFloatAtomicsFp64GenericMinMax : FunctionExtension<"cl_ext_float_atomics cl_khr_int64_base_atomics cl_khr_int64_extended_atomics __opencl_c_ext_fp64_local_atomic_min_max __opencl_c_ext_fp64_global_atomic_min_max">;
// Not a real extension, but a workaround to add C++ for OpenCL specific builtins.
def FuncExtOpenCLCxx : FunctionExtension<"__cplusplus">;
Index: clang/lib/Headers/opencl-c.h
===================================================================
--- clang/lib/Headers/opencl-c.h
+++ clang/lib/Headers/opencl-c.h
@@ -13832,6 +13832,7 @@
#endif // defined(__opencl_c_ext_fp32_global_atomic_min_max) && \
defined(__opencl_c_ext_fp32_local_atomic_min_max)
+#if defined(cl_khr_int64_base_atomics) && defined(cl_khr_int64_extended_atomics)
#if defined(__opencl_c_ext_fp64_global_atomic_min_max)
double __ovld atomic_fetch_min(volatile __global atomic_double *object,
double operand);
@@ -13882,6 +13883,8 @@
memory_scope scope);
#endif // defined(__opencl_c_ext_fp64_global_atomic_min_max) && \
defined(__opencl_c_ext_fp64_local_atomic_min_max)
+#endif // defined(cl_khr_int64_base_atomics) &&
+ // defined(cl_khr_int64_extended_atomics)
#if defined(__opencl_c_ext_fp16_global_atomic_add)
half __ovld atomic_fetch_add(volatile __global atomic_half *object,
@@ -13985,6 +13988,7 @@
#endif // defined(__opencl_c_ext_fp32_global_atomic_add) && \
defined(__opencl_c_ext_fp32_local_atomic_add)
+#if defined(cl_khr_int64_base_atomics) && defined(cl_khr_int64_extended_atomics)
#if defined(__opencl_c_ext_fp64_global_atomic_add)
double __ovld atomic_fetch_add(volatile __global atomic_double *object,
double operand);
@@ -14035,6 +14039,8 @@
memory_scope scope);
#endif // defined(__opencl_c_ext_fp64_global_atomic_add) && \
defined(__opencl_c_ext_fp64_local_atomic_add)
+#endif // defined(cl_khr_int64_base_atomics) &&
+ // defined(cl_khr_int64_extended_atomics)
#endif // cl_ext_float_atomics
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
