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">;
}
}
-let MaxVersion = CL20 in {
+let Extension = FuncExtOpenCLCNamedAddressSpaceBuiltins in {
defm : MathWithPointer<[GlobalAS, LocalAS, PrivateAS]>;
}
-let MinVersion = CL20 in {
+let Extension = FuncExtOpenCLCGenericAddressSpace in {
defm : MathWithPointer<[GenericAS]>;
}
}
}
-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.
}
}
-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.
-// 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
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]] =
// 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) {