[PATCH] D119398: [OpenCL] Guard atomic_double with cl_khr_int64_base_atomics and cl_khr_int64_extended_atomics
Yang Haonan via Phabricator via cfe-commits
cfe-commits at lists.llvm.org
Wed Feb 9 18:39:03 PST 2022
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
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D119398.407362.patch
Type: text/x-patch
Size: 6040 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20220210/bb2848c6/attachment.bin>
More information about the cfe-commits
mailing list