[clang] 8e60992 - [OpenCL] Make generic addrspace optional for -fdeclare-opencl-builtins

[Sven van Haastregt via cfe-commits]

Author: Sven van Haastregt
Date: 2022-01-31T10:21:05Z
New Revision: 8e6099291dcb49b90e59a591ec24c77f348239b6

URL: https://github.com/llvm/llvm-project/commit/8e6099291dcb49b90e59a591ec24c77f348239b6
DIFF: https://github.com/llvm/llvm-project/commit/8e6099291dcb49b90e59a591ec24c77f348239b6.diff

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

Currently, -fdeclare-opencl-builtins always adds the generic address
space overloads of e.g. the vload builtin functions in OpenCL 3.0
mode, even when the generic address space feature is disabled.

Guard the generic address space overloads by the
`__opencl_c_generic_address_space` feature instead of by OpenCL
version.

Guard the private, global, and local overloads using the internal
`__opencl_c_named_address_space_builtins` feature.

Differential Revision: https://reviews.llvm.org/D107769

Added: 
    

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

Removed: 
    


################################################################################
diff  --git a/clang/lib/Sema/OpenCLBuiltins.td b/clang/lib/Sema/OpenCLBuiltins.td
index df2f206041c10..cd704ba2df13d 100644
--- a/clang/lib/Sema/OpenCLBuiltins.td
+++ b/clang/lib/Sema/OpenCLBuiltins.td
@@ -85,6 +85,8 @@ def FuncExtKhrMipmapImageWrites          : FunctionExtension<"cl_khr_mipmap_imag
 def FuncExtKhrGlMsaaSharing              : FunctionExtension<"cl_khr_gl_msaa_sharing">;
 def FuncExtKhrGlMsaaSharingReadWrite     : FunctionExtension<"cl_khr_gl_msaa_sharing __opencl_c_read_write_images">;
 
+def FuncExtOpenCLCGenericAddressSpace    : FunctionExtension<"__opencl_c_generic_address_space">;
+def FuncExtOpenCLCNamedAddressSpaceBuiltins : FunctionExtension<"__opencl_c_named_address_space_builtins">;
 def FuncExtOpenCLCPipes                  : FunctionExtension<"__opencl_c_pipes">;
 def FuncExtOpenCLCWGCollectiveFunctions  : FunctionExtension<"__opencl_c_work_group_collective_functions">;
 def FuncExtOpenCLCReadWriteImages        : FunctionExtension<"__opencl_c_read_write_images">;
@@ -591,10 +593,10 @@ multiclass MathWithPointer<list<AddressSpace> addrspaces> {
   }
 }
 
-let MaxVersion = CL20 in {
+let Extension = FuncExtOpenCLCNamedAddressSpaceBuiltins in {
   defm : MathWithPointer<[GlobalAS, LocalAS, PrivateAS]>;
 }
-let MinVersion = CL20 in {
+let Extension = FuncExtOpenCLCGenericAddressSpace in {
   defm : MathWithPointer<[GenericAS]>;
 }
 
@@ -840,10 +842,10 @@ multiclass VloadVstore<list<AddressSpace> addrspaces, bit defStores> {
   }
 }
 
-let MaxVersion = CL20 in {
+let Extension = FuncExtOpenCLCNamedAddressSpaceBuiltins in {
   defm : VloadVstore<[GlobalAS, LocalAS, PrivateAS], 1>;
 }
-let MinVersion = CL20 in {
+let Extension = FuncExtOpenCLCGenericAddressSpace in {
   defm : VloadVstore<[GenericAS], 1>;
 }
 // vload with constant address space is available regardless of version.
@@ -874,10 +876,10 @@ multiclass VloadVstoreHalf<list<AddressSpace> addrspaces, bit defStores> {
   }
 }
 
-let MaxVersion = CL20 in {
+let Extension = FuncExtOpenCLCNamedAddressSpaceBuiltins in {
   defm : VloadVstoreHalf<[GlobalAS, LocalAS, PrivateAS], 1>;
 }
-let MinVersion = CL20 in {
+let Extension = FuncExtOpenCLCGenericAddressSpace in {
   defm : VloadVstoreHalf<[GenericAS], 1>;
 }
 // vload_half and vloada_half with constant address space are available regardless of version.

diff  --git a/clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl b/clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl
index 665297b469f83..f0ebf6e3c0eac 100644
--- a/clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl
+++ b/clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl
@@ -1,5 +1,12 @@
-// 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 \
+// RUN: -cl-ext=-__opencl_c_generic_address_space,-__opencl_c_pipes,-__opencl_c_device_enqueue %s \
+// RUN: | FileCheck %s --check-prefixes CHECK,CHECK-NOGAS
 
 // Test that mix is correctly defined.
 // CHECK-LABEL: @test_float
@@ -32,6 +39,15 @@ kernel void test_mangling() {
   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]] =

diff  --git a/clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl b/clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl
index be6e53a07bdf3..8c5863987be68 100644
--- a/clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl
+++ b/clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl
@@ -70,12 +70,15 @@ typedef struct {int a;} ndrange_t;
 
 // 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
 #define cl_khr_subgroup_clustered_reduce 1
 #define __opencl_c_read_write_images 1
 #endif
+
+#define __opencl_c_named_address_space_builtins 1
 #endif
 
 kernel void test_pointers(volatile global void *global_p, global const int4 *a) {