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

[svenvh]

Summary

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.

[svenvh]

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?

[Anastasia]

In D107769#2960441, @svenvh wrote:

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?

Yes, we could do this indeed as a clang extension. I don't think this will impact the user code. However, this means vendors will have to add definitions for extra overloads in OpenCL 2.0 mode?

[svenvh]

In D107769#2967565, @Anastasia wrote:

In D107769#2960441, @svenvh wrote:

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?

Yes, we could do this indeed as a clang extension. I don't think this will impact the user code. However, this means vendors will have to add definitions for extra overloads in OpenCL 2.0 mode?

I wonder if making the private, global, and local overloads always available would even be considered an extension. Unless I overlooked something, I cannot find anything in the spec saying that the private, global, and local overloads should not be available when a generic overload is available (even though this is what Clang has always done)?

Making all overloads always available in e.g. CL2.0 mode will indeed affect the generated calls, which means the definitions should be available when consuming the generated IR.

[Anastasia]

In D107769#3064171, @svenvh wrote:

In D107769#2967565, @Anastasia wrote:

In D107769#2960441, @svenvh wrote:

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?

Yes, we could do this indeed as a clang extension. I don't think this will impact the user code. However, this means vendors will have to add definitions for extra overloads in OpenCL 2.0 mode?

I wonder if making the private, global, and local overloads always available would even be considered an extension. Unless I overlooked something, I cannot find anything in the spec saying that the private, global, and local overloads should not be available when a generic overload is available (even though this is what Clang has always done)?

Making all overloads always available in e.g. CL2.0 mode will indeed affect the generated calls, which means the definitions should be available when consuming the generated IR.

The way I understand the spec for OpenCL C 2.0 is that whenever the address space of the pointer is not listed it means generic has to be used, here is one example:
https://www.khronos.org/registry/OpenCL/specs/2.2/html/OpenCL_C.html#vector-data-load-and-store-functions

gentypen vloadn(size_t offset, const gentype *p)
gentypen vloadn(size_t offset, const __constant gentype *p)

that has no address space (i.e. generic) and constant explicitly. So I think if address spaces are not listed explicitly they are not supposed to be available.

One implication of adding all address space overloads is that it makes library size larger, but my feeling is that we don't have that many functions with pointers to significantly impace the library size?

[svenvh]

In D107769#3265665, @Anastasia wrote:

The way I understand the spec for OpenCL C 2.0 is that whenever the address space of the pointer is not listed it means generic has to be used, here is one example:
https://www.khronos.org/registry/OpenCL/specs/2.2/html/OpenCL_C.html#vector-data-load-and-store-functions

gentypen vloadn(size_t offset, const gentype *p)
gentypen vloadn(size_t offset, const __constant gentype *p)

that has no address space (i.e. generic) and constant explicitly. So I think if address spaces are not listed explicitly they are not supposed to be available.

The unified specification (which "specifies all versions of OpenCL C") seems to be making all overloads available as I understand; it is perhaps subtly different from the previous specification?

https://www.khronos.org/registry/OpenCL/specs/3.0-unified/html/OpenCL_C.html#vector-data-load-and-store-functions

gentypen vloadn(size_t offset, const __global gentype *p)
gentypen vloadn(size_t offset, const __local gentype *p)
gentypen vloadn(size_t offset, const __constant gentype *p)
gentypen vloadn(size_t offset, const __private gentype *p)

For OpenCL C 2.0, or OpenCL C 3.0 or newer with the __opencl_c_generic_address_space feature:

gentypen vloadn(size_t offset, const gentype *p)

Since the __constant overload should always be available, I think a reader can assume that the overloads directly above and below __constant are also always available? So that the generic overload is an optional addition to the list of overloads. If not, I'd expect the spec to specify a condition before listing the specific overloads.

One implication of adding all address space overloads is that it makes library size larger, but my feeling is that we don't have that many functions with pointers to significantly impace the library size?

This patch should be touching all of them. Not that many indeed, but it might still have a non-negligible impact on OpenCL libraries due to the combination of #vector-sizes * #types * #addrspaces.

[Anastasia]

In D107769#3265867, @svenvh wrote:

In D107769#3265665, @Anastasia wrote:

The way I understand the spec for OpenCL C 2.0 is that whenever the address space of the pointer is not listed it means generic has to be used, here is one example:
https://www.khronos.org/registry/OpenCL/specs/2.2/html/OpenCL_C.html#vector-data-load-and-store-functions

gentypen vloadn(size_t offset, const gentype *p)
gentypen vloadn(size_t offset, const __constant gentype *p)

that has no address space (i.e. generic) and constant explicitly. So I think if address spaces are not listed explicitly they are not supposed to be available.

The unified specification (which "specifies all versions of OpenCL C") seems to be making all overloads available as I understand; it is perhaps subtly different from the previous specification?

https://www.khronos.org/registry/OpenCL/specs/3.0-unified/html/OpenCL_C.html#vector-data-load-and-store-functions

gentypen vloadn(size_t offset, const __global gentype *p)
gentypen vloadn(size_t offset, const __local gentype *p)
gentypen vloadn(size_t offset, const __constant gentype *p)
gentypen vloadn(size_t offset, const __private gentype *p)

For OpenCL C 2.0, or OpenCL C 3.0 or newer with the __opencl_c_generic_address_space feature:

gentypen vloadn(size_t offset, const gentype *p)

Since the __constant overload should always be available, I think a reader can assume that the overloads directly above and below __constant are also always available? So that the generic overload is an optional addition to the list of overloads. If not, I'd expect the spec to specify a condition before listing the specific overloads.

I agree that OpenCL 3.0 spec might have broken the backward compatibility for OpenCL 2.0, however, I am not convinced it has been done intentionally. Perhaps, it's worth to clarify that?

[Anastasia]

Looping in @yaxunl for further feedback regarding OpenCL 2.0 behavior specifically.

[svenvh]

Make use of the __opencl_c_named_address_space_builtins internal feature added by D118158. This should avoid affecting OpenCL 2.0.

[Anastasia]

LGTM. Thanks for changing this. I have no issues with this patch as it doesn't change the behavior for OpenCL 2.0 but I did create an issue for OpenCL-Docs to fix this problem somehow: https://github.com/KhronosGroup/OpenCL-Docs/issues/753. So in case, there is any feedback we can refine this behavior later on.

clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl

73

btw this is not correct for C++ for OpenCL 2021 but we are not testing this with C++ for OpenCL 2021 which we should.

However it doesn't belong to this patch, but would you be able to add a FIXME here to indicate the issue?

[svenvh]

clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl

73

I'll add C++ for OpenCL 2021 (and CL3.0) testing in a separate commit.

As discussed offline, we are likely to drop headerless testing from here soonish, since we're relying more and more on opencl-c-base.h for -fdeclare-opencl-builtins. We essentially have to redefine all the feature macros in the test too, when they're already in the header that we're deliberately excluding.

clang/lib/Sema/OpenCLBuiltins.td

 def FuncExtKhrInt64BaseAtomics           : FunctionExtension<"cl_khr_int64_base_atomics">;
 def FuncExtKhrInt64ExtendedAtomics       : FunctionExtension<"cl_khr_int64_extended_atomics">;
 def FuncExtKhrMipmapImage                : FunctionExtension<"cl_khr_mipmap_image">;
 def FuncExtKhrMipmapImageReadWrite       : FunctionExtension<"cl_khr_mipmap_image __opencl_c_read_write_images">;
 def FuncExtKhrMipmapImageWrites          : FunctionExtension<"cl_khr_mipmap_image_writes">;
 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">;
 def FuncExtFloatAtomicsFp16GlobalLoadStore  : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_global_atomic_load_store">;
 def FuncExtFloatAtomicsFp16LocalLoadStore   : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_local_atomic_load_store">;
 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">;
foreach AS = addrspaces in {
     foreach name = ["remquo"] in {
       foreach Type = [GenTypeFloatVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeHalfVecAndScalar] in {
         def : Builtin<name, [Type, Type, Type, PointerType<GenTypeIntVecAndScalar, AS>]>;
       }
     }
   }
 }
 
-let MaxVersion = CL20 in {
+let Extension = FuncExtOpenCLCNamedAddressSpaceBuiltins in {
   defm : MathWithPointer<[GlobalAS, LocalAS, PrivateAS]>;
 }
-let MinVersion = CL20 in {
+let Extension = FuncExtOpenCLCGenericAddressSpace in {
   defm : MathWithPointer<[GenericAS]>;
 }
 
 // --- Table 9 ---
 foreach name = ["half_cos",
                 "half_exp", "half_exp2", "half_exp10",
                 "half_log", "half_log2", "half_log10",
                 "half_recip", "half_rsqrt",
foreach VSize = [2, 3, 4, 8, 16] in {
           def : Builtin<name, [Void, VectorType<Double, VSize>, Size, PointerType<Double, AS>]>;
           def : Builtin<name, [Void, VectorType<Half, VSize>, Size, PointerType<Half, AS>]>;
         }
       }
     }
   }
 }
 
-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.
 defm : VloadVstore<[ConstantAS], 0>;
 
 multiclass VloadVstoreHalf<list<AddressSpace> addrspaces, bit defStores> {
   foreach AS = addrspaces in {
     def : Builtin<"vload_half", [Float, Size, PointerType<ConstType<Half>, AS>], Attr.Pure>;
if defStores then {
             def : Builtin<name, [Void, VectorType<Double, VSize>, Size, PointerType<Half, AS>]>;
           }
         }
       }
     }
   }
 }
 
-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.
 defm : VloadVstoreHalf<[ConstantAS], 0>;
 
 // OpenCL v3.0 s6.15.8 - Synchronization Functions.
 def : Builtin<"barrier", [Void, MemFenceFlags], Attr.Convergent>;
 let MinVersion = CL20 in {

clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl

-// 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
 // CHECK: call spir_func <4 x float> @_Z3mixDv4_fS_f
 // CHECK: ret
 void test_float(float4 x, float a) {
   float4 ret = mix(x, x, a);
 }

 
 // Test that builtins with only one prototype are mangled.
 // CHECK-LABEL: @test_mangling
 // CHECK: call spir_func i32 @_Z12get_local_idj
 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]] =
 // CHECK-SAME: readonly

clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl

 
 typedef uint cl_mem_fence_flags;
 #define CLK_GLOBAL_MEM_FENCE 0x02
 
 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) {
   int i;
   unsigned int ui;
 
   prefetch(a, 2);