[PATCH] D107769: [OpenCL] Make generic addrspace optional for -fdeclare-opencl-builtins

[Sven van Haastregt via Phabricator via cfe-commits]

svenvh updated this revision to Diff 368136.
svenvh edited the summary of this revision.
svenvh added a comment.

I have done an alternative spin of this patch: instead of introducing `RequireDisabledExtension`, simply always make the `private`, `global`, and `local` overloads available.  This makes tablegen diverge from `opencl-c.h` though.  Perhaps we should also always add make the `private`, `global`, and `local` overloads available in `opencl-c.h`?


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D107769/new/

https://reviews.llvm.org/D107769

Files:
  clang/lib/Sema/OpenCLBuiltins.td
  clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl
  clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl


Index: clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl
===================================================================
--- clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl
+++ clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl
@@ -63,6 +63,7 @@
 
 // Enable extensions that are enabled in opencl-c-base.h.
 #if (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200)
+#define __opencl_c_generic_address_space 1
 #define cl_khr_subgroup_extended_types 1
 #define cl_khr_subgroup_ballot 1
 #define cl_khr_subgroup_non_uniform_arithmetic 1
Index: clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl
===================================================================
--- clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl
+++ clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl
@@ -1,5 +1,11 @@
-// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL1.2 -finclude-default-header %s | FileCheck %s
-// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL1.2 -fdeclare-opencl-builtins -finclude-default-header %s | FileCheck %s
+// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL1.2 -finclude-default-header %s \
+// RUN: | FileCheck %s --check-prefixes CHECK,CHECK-NOGAS
+// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL1.2 -fdeclare-opencl-builtins -finclude-default-header %s \
+// RUN: | FileCheck %s --check-prefixes CHECK,CHECK-NOGAS
+// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL3.0 -fdeclare-opencl-builtins -finclude-default-header %s \
+// RUN: | FileCheck %s --check-prefixes CHECK,CHECK-GAS
+// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL3.0 -fdeclare-opencl-builtins -finclude-default-header -cl-ext=-__opencl_c_generic_address_space,-__opencl_c_pipes %s \
+// RUN: | FileCheck %s --check-prefixes CHECK,CHECK-NOGAS
 
 // Test that mix is correctly defined.
 // CHECK-LABEL: @test_float
@@ -32,6 +38,15 @@
   size_t lid = get_local_id(0);
 }
 
+// Test that the correct builtin is called depending on the generic address
+// space feature availability.
+// CHECK-LABEL: @test_generic_optionality
+// CHECK-GAS: call spir_func float @_Z5fractfPU3AS4f
+// CHECK-NOGAS: call spir_func float @_Z5fractfPf
+void test_generic_optionality(float a, float *b) {
+  float res = fract(a, b);
+}
+
 // CHECK: attributes [[ATTR_CONST]] =
 // CHECK-SAME: readnone
 // CHECK: attributes [[ATTR_PURE]] =
Index: clang/lib/Sema/OpenCLBuiltins.td
===================================================================
--- clang/lib/Sema/OpenCLBuiltins.td
+++ clang/lib/Sema/OpenCLBuiltins.td
@@ -83,6 +83,7 @@
 def FuncExtKhrMipmapImageWrites          : FunctionExtension<"cl_khr_mipmap_image_writes">;
 def FuncExtKhrGlMsaaSharing              : FunctionExtension<"cl_khr_gl_msaa_sharing">;
 
+def FuncExtOpenCLCGenericAddressSpace    : FunctionExtension<"__opencl_c_generic_address_space">;
 def FuncExtOpenCLCPipes                  : FunctionExtension<"__opencl_c_pipes">;
 def FuncExtOpenCLCWGCollectiveFunctions  : FunctionExtension<"__opencl_c_work_group_collective_functions">;
 
@@ -566,10 +567,8 @@
   }
 }
 
-let MaxVersion = CL20 in {
-  defm : MathWithPointer<[GlobalAS, LocalAS, PrivateAS]>;
-}
-let MinVersion = CL20 in {
+defm : MathWithPointer<[GlobalAS, LocalAS, PrivateAS]>;
+let Extension = FuncExtOpenCLCGenericAddressSpace in {
   defm : MathWithPointer<[GenericAS]>;
 }
 
@@ -824,10 +823,8 @@
   }
 }
 
-let MaxVersion = CL20 in {
-  defm : VloadVstore<[GlobalAS, LocalAS, PrivateAS], 1>;
-}
-let MinVersion = CL20 in {
+defm : VloadVstore<[GlobalAS, LocalAS, PrivateAS], 1>;
+let Extension = FuncExtOpenCLCGenericAddressSpace in {
   defm : VloadVstore<[GenericAS], 1>;
 }
 // vload with constant address space is available regardless of version.
@@ -859,10 +856,8 @@
   }
 }
 
-let MaxVersion = CL20 in {
-  defm : VloadVstoreHalf<[GlobalAS, LocalAS, PrivateAS], 1>;
-}
-let MinVersion = CL20 in {
+defm : VloadVstoreHalf<[GlobalAS, LocalAS, PrivateAS], 1>;
+let Extension = FuncExtOpenCLCGenericAddressSpace in {
   defm : VloadVstoreHalf<[GenericAS], 1>;
 }
 // vload with constant address space is available regardless of version.


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D107769.368136.patch
Type: text/x-patch
Size: 4233 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20210823/65ea9334/attachment-0001.bin>