64 lines · plain
1// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL1.2 -finclude-default-header %s \2// RUN: | FileCheck %s --check-prefixes CHECK,CHECK-NOGAS3// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL1.2 -fdeclare-opencl-builtins -finclude-default-header %s \4// RUN: | FileCheck %s --check-prefixes CHECK,CHECK-NOGAS5// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL3.0 -fdeclare-opencl-builtins -finclude-default-header %s \6// RUN: | FileCheck %s --check-prefixes CHECK,CHECK-GAS7// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL3.0 -fdeclare-opencl-builtins -finclude-default-header \8// RUN: -cl-ext=-__opencl_c_generic_address_space,-__opencl_c_pipes,-__opencl_c_device_enqueue %s \9// RUN: | FileCheck %s --check-prefixes CHECK,CHECK-NOGAS10 11// Test that mix is correctly defined.12// CHECK-LABEL: @test_float13// CHECK: call spir_func <4 x float> @_Z3mixDv4_fS_f14// CHECK: ret15void test_float(float4 x, float a) {16 float4 ret = mix(x, x, a);17}18 19// Test that Attr.Const from OpenCLBuiltins.td is lowered to a readnone attribute.20// CHECK-LABEL: @test_const_attr21// CHECK: call spir_func i32 @_Z3maxii({{.*}}) [[ATTR_CONST:#[0-9]]]22// CHECK: ret23int test_const_attr(int a) {24 return max(a, 2);25}26 27// Test that Attr.Pure from OpenCLBuiltins.td is lowered to a readonly attribute.28// CHECK-LABEL: @test_pure_attr29// CHECK: call spir_func <4 x float> @_Z11read_imagef{{.*}} [[ATTR_PURE:#[0-9]]]30// CHECK: ret31kernel void test_pure_attr(read_only image1d_t img) {32 float4 resf = read_imagef(img, 42);33}34 35// Test that builtins with only one prototype are mangled.36// CHECK-LABEL: @test_mangling37// CHECK: call spir_func i32 @_Z12get_local_idj38kernel void test_mangling() {39 size_t lid = get_local_id(0);40}41 42// Test that the correct builtin is called depending on the generic address43// space feature availability.44// CHECK-LABEL: @test_generic_optionality45// CHECK-GAS: call spir_func float @_Z5fractfPU3AS4f46// CHECK-NOGAS: call spir_func float @_Z5fractfPf47void test_generic_optionality(float a, float *b) {48 float res = fract(a, b);49}50 51// Test that the correct builtin is called depending on the generic address52// space feature availability. If not available, the __private version is called53// CHECK-LABEL: @test_wait_group_events54// CHECK-GAS: call spir_func void @_Z17wait_group_eventsiPU3AS49ocl_event55// CHECK-NOGAS: call spir_func void @_Z17wait_group_eventsiP9ocl_event56void test_wait_group_events(int i, event_t *e) {57 wait_group_events(i, e);58}59 60// CHECK: attributes [[ATTR_CONST]] =61// CHECK-SAME: memory(none)62// CHECK: attributes [[ATTR_PURE]] =63// CHECK-SAME: memory(read)64