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.
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.
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
