diff options
Diffstat (limited to 'test/CodeGenOpenCL')
34 files changed, 1673 insertions, 210 deletions
diff --git a/test/CodeGenOpenCL/address-space-constant-initializers.cl b/test/CodeGenOpenCL/address-space-constant-initializers.cl index 44c4ddf878005..45c64c2f716b2 100644 --- a/test/CodeGenOpenCL/address-space-constant-initializers.cl +++ b/test/CodeGenOpenCL/address-space-constant-initializers.cl @@ -11,8 +11,8 @@ typedef struct { __constant float* constant_float_ptr; } ConstantArrayPointerStruct; -// CHECK: %struct.ConstantArrayPointerStruct = type { float addrspace(3)* } -// CHECK: addrspace(3) constant %struct.ConstantArrayPointerStruct { float addrspace(3)* bitcast (i8 addrspace(3)* getelementptr (i8, i8 addrspace(3)* bitcast (%struct.ArrayStruct addrspace(3)* @constant_array_struct to i8 addrspace(3)*), i64 4) to float addrspace(3)*) } +// CHECK: %struct.ConstantArrayPointerStruct = type { float addrspace(2)* } +// CHECK: addrspace(2) constant %struct.ConstantArrayPointerStruct { float addrspace(2)* bitcast (i8 addrspace(2)* getelementptr (i8, i8 addrspace(2)* bitcast (%struct.ArrayStruct addrspace(2)* @constant_array_struct to i8 addrspace(2)*), i64 4) to float addrspace(2)*) } // Bug 18567 __constant ConstantArrayPointerStruct constant_array_pointer_struct = { &constant_array_struct.f diff --git a/test/CodeGenOpenCL/address-spaces-mangling.cl b/test/CodeGenOpenCL/address-spaces-mangling.cl index edb53fc35821e..3c74c718c2a2a 100644 --- a/test/CodeGenOpenCL/address-spaces-mangling.cl +++ b/test/CodeGenOpenCL/address-spaces-mangling.cl @@ -1,6 +1,10 @@ // RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=yes -triple %itanium_abi_triple -emit-llvm -o - | FileCheck -check-prefix=ASMANG %s // RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=no -triple %itanium_abi_triple -emit-llvm -o - | FileCheck -check-prefix=NOASMANG %s +// We check that the address spaces are mangled the same in both version of OpenCL +// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=CL2.0 -emit-llvm -o - | FileCheck -check-prefix=OCL-20 %s +// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=CL1.2 -emit-llvm -o - | FileCheck -check-prefix=OCL-12 %s + // We can't name this f as private is equivalent to default // no specifier given address space so we get multiple definition // warnings, but we do want it for comparison purposes. @@ -8,23 +12,33 @@ __attribute__((overloadable)) void ff(int *arg) { } // ASMANG: @_Z2ffPi // NOASMANG: @_Z2ffPi +// OCL-20-DAG: @_Z2ffPU3AS4i +// OCL-12-DAG: @_Z2ffPi __attribute__((overloadable)) void f(private int *arg) { } // ASMANG: @_Z1fPi // NOASMANG: @_Z1fPi +// OCL-20-DAG: @_Z1fPi +// OCL-12-DAG: @_Z1fPi __attribute__((overloadable)) void f(global int *arg) { } // ASMANG: @_Z1fPU3AS1i // NOASMANG: @_Z1fPU8CLglobali +// OCL-20-DAG: @_Z1fPU3AS1i +// OCL-12-DAG: @_Z1fPU3AS1i __attribute__((overloadable)) void f(local int *arg) { } -// ASMANG: @_Z1fPU3AS2i +// ASMANG: @_Z1fPU3AS3i // NOASMANG: @_Z1fPU7CLlocali +// OCL-20-DAG: @_Z1fPU3AS3i +// OCL-12-DAG: @_Z1fPU3AS3i __attribute__((overloadable)) void f(constant int *arg) { } -// ASMANG: @_Z1fPU3AS3i +// ASMANG: @_Z1fPU3AS2i // NOASMANG: @_Z1fPU10CLconstanti +// OCL-20-DAG: @_Z1fPU3AS2i +// OCL-12-DAG: @_Z1fPU3AS2i diff --git a/test/CodeGenOpenCL/address-spaces.cl b/test/CodeGenOpenCL/address-spaces.cl index 68fa02dadb26d..3b34986d89360 100644 --- a/test/CodeGenOpenCL/address-spaces.cl +++ b/test/CodeGenOpenCL/address-spaces.cl @@ -7,10 +7,10 @@ void f__p(__private int *arg) {} // CHECK: i32 addrspace(1)* %arg void f__g(__global int *arg) {} -// CHECK: i32 addrspace(2)* %arg +// CHECK: i32 addrspace(3)* %arg void f__l(__local int *arg) {} -// CHECK: i32 addrspace(3)* %arg +// CHECK: i32 addrspace(2)* %arg void f__c(__constant int *arg) {} // CHECK: i32* %arg @@ -19,10 +19,10 @@ void fp(private int *arg) {} // CHECK: i32 addrspace(1)* %arg void fg(global int *arg) {} -// CHECK: i32 addrspace(2)* %arg +// CHECK: i32 addrspace(3)* %arg void fl(local int *arg) {} -// CHECK: i32 addrspace(3)* %arg +// CHECK: i32 addrspace(2)* %arg void fc(constant int *arg) {} #ifdef CL20 diff --git a/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl b/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl new file mode 100644 index 0000000000000..f37fe66239920 --- /dev/null +++ b/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl @@ -0,0 +1,66 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s + +// CHECK-NOT: %struct.single_element_struct_arg = type { i32 } +typedef struct single_element_struct_arg +{ + int i; +} single_element_struct_arg_t; + +// CHECK: %struct.struct_arg = type { i32, float, i32 } +typedef struct struct_arg +{ + int i1; + float f; + int i2; +} struct_arg_t; + +// CHECK: %struct.struct_of_arrays_arg = type { [2 x i32], float, [4 x i32], [3 x float], i32 } +typedef struct struct_of_arrays_arg +{ + int i1[2]; + float f1; + int i2[4]; + float f2[3]; + int i3; +} struct_of_arrays_arg_t; + +// CHECK: %struct.struct_of_structs_arg = type { i32, float, %struct.struct_arg, i32 } +typedef struct struct_of_structs_arg +{ + int i1; + float f1; + struct_arg_t s1; + int i2; +} struct_of_structs_arg_t; + +// CHECK-LABEL: @test_single_element_struct_arg +// CHECK: i32 %arg1.coerce +__kernel void test_single_element_struct_arg(single_element_struct_arg_t arg1) +{ +} + +// CHECK-LABEL: @test_struct_arg +// CHECK: %struct.struct_arg %arg1.coerce +__kernel void test_struct_arg(struct_arg_t arg1) +{ +} + +// CHECK-LABEL: @test_struct_of_arrays_arg +// CHECK: %struct.struct_of_arrays_arg %arg1.coerce +__kernel void test_struct_of_arrays_arg(struct_of_arrays_arg_t arg1) +{ +} + +// CHECK-LABEL: @test_struct_of_structs_arg +// CHECK: %struct.struct_of_structs_arg %arg1.coerce +__kernel void test_struct_of_structs_arg(struct_of_structs_arg_t arg1) +{ +} + +// CHECK-LABEL: @test_non_kernel_struct_arg +// CHECK-NOT: %struct.struct_arg %arg1.coerce +// CHECK: %struct.struct_arg* byval +void test_non_kernel_struct_arg(struct_arg_t arg1) +{ +} diff --git a/test/CodeGenOpenCL/amdgpu-attrs.cl b/test/CodeGenOpenCL/amdgpu-attrs.cl new file mode 100644 index 0000000000000..4ca85d316104e --- /dev/null +++ b/test/CodeGenOpenCL/amdgpu-attrs.cl @@ -0,0 +1,166 @@ +// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -O0 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -verify -o - %s | FileCheck -check-prefix=X86 %s + +__attribute__((amdgpu_flat_work_group_size(0, 0))) // expected-no-diagnostics +kernel void flat_work_group_size_0_0() {} +__attribute__((amdgpu_waves_per_eu(0))) // expected-no-diagnostics +kernel void waves_per_eu_0() {} +__attribute__((amdgpu_waves_per_eu(0, 0))) // expected-no-diagnostics +kernel void waves_per_eu_0_0() {} +__attribute__((amdgpu_num_sgpr(0))) // expected-no-diagnostics +kernel void num_sgpr0() {} +__attribute__((amdgpu_num_vgpr(0))) // expected-no-diagnostics +kernel void num_vgpr0() {} + +__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0))) // expected-no-diagnostics +kernel void flat_work_group_size_0_0_waves_per_eu_0() {} +__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0, 0))) // expected-no-diagnostics +kernel void flat_work_group_size_0_0_waves_per_eu_0_0() {} +__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_num_sgpr(0))) // expected-no-diagnostics +kernel void flat_work_group_size_0_0_num_sgpr_0() {} +__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_num_vgpr(0))) // expected-no-diagnostics +kernel void flat_work_group_size_0_0_num_vgpr_0() {} +__attribute__((amdgpu_waves_per_eu(0), amdgpu_num_sgpr(0))) // expected-no-diagnostics +kernel void waves_per_eu_0_num_sgpr_0() {} +__attribute__((amdgpu_waves_per_eu(0), amdgpu_num_vgpr(0))) // expected-no-diagnostics +kernel void waves_per_eu_0_num_vgpr_0() {} +__attribute__((amdgpu_waves_per_eu(0, 0), amdgpu_num_sgpr(0))) // expected-no-diagnostics +kernel void waves_per_eu_0_0_num_sgpr_0() {} +__attribute__((amdgpu_waves_per_eu(0, 0), amdgpu_num_vgpr(0))) // expected-no-diagnostics +kernel void waves_per_eu_0_0_num_vgpr_0() {} +__attribute__((amdgpu_num_sgpr(0), amdgpu_num_vgpr(0))) // expected-no-diagnostics +kernel void num_sgpr_0_num_vgpr_0() {} + +__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0), amdgpu_num_sgpr(0))) // expected-no-diagnostics +kernel void flat_work_group_size_0_0_waves_per_eu_0_num_sgpr_0() {} +__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0), amdgpu_num_vgpr(0))) // expected-no-diagnostics +kernel void flat_work_group_size_0_0_waves_per_eu_0_num_vgpr_0() {} +__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0, 0), amdgpu_num_sgpr(0))) // expected-no-diagnostics +kernel void flat_work_group_size_0_0_waves_per_eu_0_0_num_sgpr_0() {} +__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0, 0), amdgpu_num_vgpr(0))) // expected-no-diagnostics +kernel void flat_work_group_size_0_0_waves_per_eu_0_0_num_vgpr_0() {} + +__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0), amdgpu_num_sgpr(0), amdgpu_num_vgpr(0))) // expected-no-diagnostics +kernel void flat_work_group_size_0_0_waves_per_eu_0_num_sgpr_0_num_vgpr_0() {} +__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0, 0), amdgpu_num_sgpr(0), amdgpu_num_vgpr(0))) // expected-no-diagnostics +kernel void flat_work_group_size_0_0_waves_per_eu_0_0_num_sgpr_0_num_vgpr_0() {} + +__attribute__((amdgpu_flat_work_group_size(32, 64))) // expected-no-diagnostics +kernel void flat_work_group_size_32_64() { +// CHECK: define amdgpu_kernel void @flat_work_group_size_32_64() [[FLAT_WORK_GROUP_SIZE_32_64:#[0-9]+]] +} +__attribute__((amdgpu_waves_per_eu(2))) // expected-no-diagnostics +kernel void waves_per_eu_2() { +// CHECK: define amdgpu_kernel void @waves_per_eu_2() [[WAVES_PER_EU_2:#[0-9]+]] +} +__attribute__((amdgpu_waves_per_eu(2, 4))) // expected-no-diagnostics +kernel void waves_per_eu_2_4() { +// CHECK: define amdgpu_kernel void @waves_per_eu_2_4() [[WAVES_PER_EU_2_4:#[0-9]+]] +} +__attribute__((amdgpu_num_sgpr(32))) // expected-no-diagnostics +kernel void num_sgpr_32() { +// CHECK: define amdgpu_kernel void @num_sgpr_32() [[NUM_SGPR_32:#[0-9]+]] +} +__attribute__((amdgpu_num_vgpr(64))) // expected-no-diagnostics +kernel void num_vgpr_64() { +// CHECK: define amdgpu_kernel void @num_vgpr_64() [[NUM_VGPR_64:#[0-9]+]] +} + +__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2))) // expected-no-diagnostics +kernel void flat_work_group_size_32_64_waves_per_eu_2() { +// CHECK: define amdgpu_kernel void @flat_work_group_size_32_64_waves_per_eu_2() [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2:#[0-9]+]] +} +__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2, 4))) // expected-no-diagnostics +kernel void flat_work_group_size_32_64_waves_per_eu_2_4() { +// CHECK: define amdgpu_kernel void @flat_work_group_size_32_64_waves_per_eu_2_4() [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4:#[0-9]+]] +} +__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_num_sgpr(32))) // expected-no-diagnostics +kernel void flat_work_group_size_32_64_num_sgpr_32() { +// CHECK: define amdgpu_kernel void @flat_work_group_size_32_64_num_sgpr_32() [[FLAT_WORK_GROUP_SIZE_32_64_NUM_SGPR_32:#[0-9]+]] +} +__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_num_vgpr(64))) // expected-no-diagnostics +kernel void flat_work_group_size_32_64_num_vgpr_64() { +// CHECK: define amdgpu_kernel void @flat_work_group_size_32_64_num_vgpr_64() [[FLAT_WORK_GROUP_SIZE_32_64_NUM_VGPR_64:#[0-9]+]] +} +__attribute__((amdgpu_waves_per_eu(2), amdgpu_num_sgpr(32))) // expected-no-diagnostics +kernel void waves_per_eu_2_num_sgpr_32() { +// CHECK: define amdgpu_kernel void @waves_per_eu_2_num_sgpr_32() [[WAVES_PER_EU_2_NUM_SGPR_32:#[0-9]+]] +} +__attribute__((amdgpu_waves_per_eu(2), amdgpu_num_vgpr(64))) // expected-no-diagnostics +kernel void waves_per_eu_2_num_vgpr_64() { +// CHECK: define amdgpu_kernel void @waves_per_eu_2_num_vgpr_64() [[WAVES_PER_EU_2_NUM_VGPR_64:#[0-9]+]] +} +__attribute__((amdgpu_waves_per_eu(2, 4), amdgpu_num_sgpr(32))) // expected-no-diagnostics +kernel void waves_per_eu_2_4_num_sgpr_32() { +// CHECK: define amdgpu_kernel void @waves_per_eu_2_4_num_sgpr_32() [[WAVES_PER_EU_2_4_NUM_SGPR_32:#[0-9]+]] +} +__attribute__((amdgpu_waves_per_eu(2, 4), amdgpu_num_vgpr(64))) // expected-no-diagnostics +kernel void waves_per_eu_2_4_num_vgpr_64() { +// CHECK: define amdgpu_kernel void @waves_per_eu_2_4_num_vgpr_64() [[WAVES_PER_EU_2_4_NUM_VGPR_64:#[0-9]+]] +} +__attribute__((amdgpu_num_sgpr(32), amdgpu_num_vgpr(64))) // expected-no-diagnostics +kernel void num_sgpr_32_num_vgpr_64() { +// CHECK: define amdgpu_kernel void @num_sgpr_32_num_vgpr_64() [[NUM_SGPR_32_NUM_VGPR_64:#[0-9]+]] +} + +__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2), amdgpu_num_sgpr(32))) +kernel void flat_work_group_size_32_64_waves_per_eu_2_num_sgpr_32() { +// CHECK: define amdgpu_kernel void @flat_work_group_size_32_64_waves_per_eu_2_num_sgpr_32() [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_NUM_SGPR_32:#[0-9]+]] +} +__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2), amdgpu_num_vgpr(64))) +kernel void flat_work_group_size_32_64_waves_per_eu_2_num_vgpr_64() { +// CHECK: define amdgpu_kernel void @flat_work_group_size_32_64_waves_per_eu_2_num_vgpr_64() [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_NUM_VGPR_64:#[0-9]+]] +} +__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2, 4), amdgpu_num_sgpr(32))) +kernel void flat_work_group_size_32_64_waves_per_eu_2_4_num_sgpr_32() { +// CHECK: define amdgpu_kernel void @flat_work_group_size_32_64_waves_per_eu_2_4_num_sgpr_32() [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4_NUM_SGPR_32:#[0-9]+]] +} +__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2, 4), amdgpu_num_vgpr(64))) +kernel void flat_work_group_size_32_64_waves_per_eu_2_4_num_vgpr_64() { +// CHECK: define amdgpu_kernel void @flat_work_group_size_32_64_waves_per_eu_2_4_num_vgpr_64() [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4_NUM_VGPR_64:#[0-9]+]] +} + +__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2), amdgpu_num_sgpr(32), amdgpu_num_vgpr(64))) // expected-no-diagnostics +kernel void flat_work_group_size_32_64_waves_per_eu_2_num_sgpr_32_num_vgpr_64() { +// CHECK: define amdgpu_kernel void @flat_work_group_size_32_64_waves_per_eu_2_num_sgpr_32_num_vgpr_64() [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_NUM_SGPR_32_NUM_VGPR_64:#[0-9]+]] +} +__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2, 4), amdgpu_num_sgpr(32), amdgpu_num_vgpr(64))) // expected-no-diagnostics +kernel void flat_work_group_size_32_64_waves_per_eu_2_4_num_sgpr_32_num_vgpr_64() { +// CHECK: define amdgpu_kernel void @flat_work_group_size_32_64_waves_per_eu_2_4_num_sgpr_32_num_vgpr_64() [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4_NUM_SGPR_32_NUM_VGPR_64:#[0-9]+]] +} + +// Make sure this is silently accepted on other targets. +// X86-NOT: "amdgpu-flat-work-group-size" +// X86-NOT: "amdgpu-waves-per-eu" +// X86-NOT: "amdgpu-num-vgpr" +// X86-NOT: "amdgpu-num-sgpr" + +// CHECK-NOT: "amdgpu-flat-work-group-size"="0,0" +// CHECK-NOT: "amdgpu-waves-per-eu"="0" +// CHECK-NOT: "amdgpu-waves-per-eu"="0,0" +// CHECK-NOT: "amdgpu-num-sgpr"="0" +// CHECK-NOT: "amdgpu-num-vgpr"="0" + +// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64]] = { noinline nounwind "amdgpu-flat-work-group-size"="32,64" +// CHECK-DAG: attributes [[WAVES_PER_EU_2]] = { noinline nounwind "amdgpu-waves-per-eu"="2" +// CHECK-DAG: attributes [[WAVES_PER_EU_2_4]] = { noinline nounwind "amdgpu-waves-per-eu"="2,4" +// CHECK-DAG: attributes [[NUM_SGPR_32]] = { noinline nounwind "amdgpu-num-sgpr"="32" +// CHECK-DAG: attributes [[NUM_VGPR_64]] = { noinline nounwind "amdgpu-num-vgpr"="64" + +// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2]] = { noinline nounwind "amdgpu-flat-work-group-size"="32,64" "amdgpu-waves-per-eu"="2" +// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4]] = { noinline nounwind "amdgpu-flat-work-group-size"="32,64" "amdgpu-waves-per-eu"="2,4" +// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_NUM_SGPR_32]] = { noinline nounwind "amdgpu-flat-work-group-size"="32,64" "amdgpu-num-sgpr"="32" +// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_NUM_VGPR_64]] = { noinline nounwind "amdgpu-flat-work-group-size"="32,64" "amdgpu-num-vgpr"="64" +// CHECK-DAG: attributes [[WAVES_PER_EU_2_NUM_SGPR_32]] = { noinline nounwind "amdgpu-num-sgpr"="32" "amdgpu-waves-per-eu"="2" +// CHECK-DAG: attributes [[WAVES_PER_EU_2_NUM_VGPR_64]] = { noinline nounwind "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2" +// CHECK-DAG: attributes [[WAVES_PER_EU_2_4_NUM_SGPR_32]] = { noinline nounwind "amdgpu-num-sgpr"="32" "amdgpu-waves-per-eu"="2,4" +// CHECK-DAG: attributes [[WAVES_PER_EU_2_4_NUM_VGPR_64]] = { noinline nounwind "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2,4" +// CHECK-DAG: attributes [[NUM_SGPR_32_NUM_VGPR_64]] = { noinline nounwind "amdgpu-num-sgpr"="32" "amdgpu-num-vgpr"="64" + +// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_NUM_SGPR_32]] = { noinline nounwind "amdgpu-flat-work-group-size"="32,64" "amdgpu-num-sgpr"="32" "amdgpu-waves-per-eu"="2" +// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_NUM_VGPR_64]] = { noinline nounwind "amdgpu-flat-work-group-size"="32,64" "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2" +// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4_NUM_SGPR_32]] = { noinline nounwind "amdgpu-flat-work-group-size"="32,64" "amdgpu-num-sgpr"="32" "amdgpu-waves-per-eu"="2,4" +// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4_NUM_VGPR_64]] = { noinline nounwind "amdgpu-flat-work-group-size"="32,64" "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2,4" + +// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_NUM_SGPR_32_NUM_VGPR_64]] = { noinline nounwind "amdgpu-flat-work-group-size"="32,64" "amdgpu-num-sgpr"="32" "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2" +// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4_NUM_SGPR_32_NUM_VGPR_64]] = { noinline nounwind "amdgpu-flat-work-group-size"="32,64" "amdgpu-num-sgpr"="32" "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2,4" diff --git a/test/CodeGenOpenCL/amdgpu-nullptr.cl b/test/CodeGenOpenCL/amdgpu-nullptr.cl new file mode 100644 index 0000000000000..2d36fa90a6ab3 --- /dev/null +++ b/test/CodeGenOpenCL/amdgpu-nullptr.cl @@ -0,0 +1,534 @@ +// RUN: %clang_cc1 %s -cl-std=CL2.0 -include opencl-c.h -triple amdgcn -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -O0 -cl-std=CL2.0 -include opencl-c.h -triple amdgcn -emit-llvm -o - | FileCheck --check-prefix=NOOPT %s + +typedef struct { + private char *p1; + local char *p2; + constant char *p3; + global char *p4; + generic char *p5; +} StructTy1; + +typedef struct { + constant char *p3; + global char *p4; + generic char *p5; +} StructTy2; + +// LLVM requests global variable with common linkage to be initialized with zeroinitializer, therefore use -fno-common +// to suppress common linkage for tentative definition. + +// Test 0 as initializer. + +// CHECK: @private_p = local_unnamed_addr addrspace(1) global i8* addrspacecast (i8 addrspace(4)* null to i8*), align 4 +private char *private_p = 0; + +// CHECK: @local_p = local_unnamed_addr addrspace(1) global i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), align 4 +local char *local_p = 0; + +// CHECK: @global_p = local_unnamed_addr addrspace(1) global i8 addrspace(1)* null, align 4 +global char *global_p = 0; + +// CHECK: @constant_p = local_unnamed_addr addrspace(1) global i8 addrspace(2)* null, align 4 +constant char *constant_p = 0; + +// CHECK: @generic_p = local_unnamed_addr addrspace(1) global i8 addrspace(4)* null, align 4 +generic char *generic_p = 0; + +// Test NULL as initializer. + +// CHECK: @private_p_NULL = local_unnamed_addr addrspace(1) global i8* addrspacecast (i8 addrspace(4)* null to i8*), align 4 +private char *private_p_NULL = NULL; + +// CHECK: @local_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), align 4 +local char *local_p_NULL = NULL; + +// CHECK: @global_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(1)* null, align 4 +global char *global_p_NULL = NULL; + +// CHECK: @constant_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(2)* null, align 4 +constant char *constant_p_NULL = NULL; + +// CHECK: @generic_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(4)* null, align 4 +generic char *generic_p_NULL = NULL; + +// Test constant folding of null pointer. +// A null pointer should be folded to a null pointer in the target address space. + +// CHECK: @fold_generic = local_unnamed_addr addrspace(1) global i32 addrspace(4)* null, align 4 +generic int *fold_generic = (global int*)(generic float*)(private char*)0; + +// CHECK: @fold_priv = local_unnamed_addr addrspace(1) global i16* addrspacecast (i16 addrspace(4)* null to i16*), align 4 +private short *fold_priv = (private short*)(generic int*)(global void*)0; + +// CHECK: @fold_priv_arith = local_unnamed_addr addrspace(1) global i8* inttoptr (i32 9 to i8*), align 4 +private char *fold_priv_arith = (private char*)0 + 10; + +// CHECK: @fold_int = local_unnamed_addr addrspace(1) global i32 13, align 4 +int fold_int = (int)(private void*)(generic char*)(global int*)0 + 14; + +// CHECK: @fold_int2 = local_unnamed_addr addrspace(1) global i32 12, align 4 +int fold_int2 = (int) ((private void*)0 + 13); + +// CHECK: @fold_int3 = local_unnamed_addr addrspace(1) global i32 -1, align 4 +int fold_int3 = (int) ((private int*)0); + +// CHECK: @fold_int4 = local_unnamed_addr addrspace(1) global i32 7, align 4 +int fold_int4 = (int) &((private int*)0)[2]; + +// CHECK: @fold_int5 = local_unnamed_addr addrspace(1) global i32 3, align 4 +int fold_int5 = (int) &((private StructTy1*)0)->p2; + +// Test static variable initialization. + +// NOOPT: @test_static_var.sp1 = internal addrspace(1) global i8* addrspacecast (i8 addrspace(4)* null to i8*), align 4 +// NOOPT: @test_static_var.sp2 = internal addrspace(1) global i8* addrspacecast (i8 addrspace(4)* null to i8*), align 4 +// NOOPT: @test_static_var.sp3 = internal addrspace(1) global i8* addrspacecast (i8 addrspace(4)* null to i8*), align 4 +// NOOPT: @test_static_var.sp4 = internal addrspace(1) global i8* null, align 4 +// NOOPT: @test_static_var.sp5 = internal addrspace(1) global i8* null, align 4 +// NOOPT: @test_static_var.SS1 = internal addrspace(1) global %struct.StructTy1 { i8* addrspacecast (i8 addrspace(4)* null to i8*), i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), i8 addrspace(2)* null, i8 addrspace(1)* null, i8 addrspace(4)* null }, align 4 +// NOOPT: @test_static_var.SS2 = internal addrspace(1) global %struct.StructTy2 zeroinitializer, align 4 + +void test_static_var(void) { + static private char *sp1 = 0; + static private char *sp2 = NULL; + static private char *sp3; + static private char *sp4 = (private char*)((void)0, 0); + const int x = 0; + static private char *sp5 = (private char*)x; + static StructTy1 SS1; + static StructTy2 SS2; +} + +// Test function-scope variable initialization. +// NOOPT-LABEL: test_func_scope_var +// NOOPT: store i8* addrspacecast (i8 addrspace(4)* null to i8*), i8** %sp1, align 4 +// NOOPT: store i8* addrspacecast (i8 addrspace(4)* null to i8*), i8** %sp2, align 4 +// NOOPT: store i8* null, i8** %sp3, align 4 +// NOOPT: store i8* null, i8** %sp4, align 4 +// NOOPT: %[[SS1:.*]] = bitcast %struct.StructTy1* %SS1 to i8* +// NOOPT: call void @llvm.memcpy.p0i8.p2i8.i64(i8* %[[SS1]], i8 addrspace(2)* bitcast (%struct.StructTy1 addrspace(2)* @test_func_scope_var.SS1 to i8 addrspace(2)*), i64 32, i32 4, i1 false) +// NOOPT: %[[SS2:.*]] = bitcast %struct.StructTy2* %SS2 to i8* +// NOOPT: call void @llvm.memset.p0i8.i64(i8* %[[SS2]], i8 0, i64 24, i32 4, i1 false) + +void test_func_scope_var(void) { + private char *sp1 = 0; + private char *sp2 = NULL; + private char *sp3 = (private char*)((void)0, 0); + const int x = 0; + private char *sp4 = (private char*)x; + StructTy1 SS1 = {0, 0, 0, 0, 0}; + StructTy2 SS2 = {0, 0, 0}; +} + +// Test default initialization of pointers. + +// Tentative definition of global variables with non-zero initializer +// cannot have common linkage since common linkage requires zero initialization +// and does not have explicit section. + +// CHECK: @p1 = weak local_unnamed_addr addrspace(1) global i8* addrspacecast (i8 addrspace(4)* null to i8*), align 4 +private char *p1; + +// CHECK: @p2 = weak local_unnamed_addr addrspace(1) global i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), align 4 +local char *p2; + +// CHECK: @p3 = common local_unnamed_addr addrspace(1) global i8 addrspace(2)* null, align 4 +constant char *p3; + +// CHECK: @p4 = common local_unnamed_addr addrspace(1) global i8 addrspace(1)* null, align 4 +global char *p4; + +// CHECK: @p5 = common local_unnamed_addr addrspace(1) global i8 addrspace(4)* null, align 4 +generic char *p5; + +// Test default initialization of sturcture. + +// CHECK: @S1 = weak local_unnamed_addr addrspace(1) global %struct.StructTy1 { i8* addrspacecast (i8 addrspace(4)* null to i8*), i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), i8 addrspace(2)* null, i8 addrspace(1)* null, i8 addrspace(4)* null }, align 4 +StructTy1 S1; + +// CHECK: @S2 = common local_unnamed_addr addrspace(1) global %struct.StructTy2 zeroinitializer, align 4 +StructTy2 S2; + +// Test default initialization of array. +// CHECK: @A1 = weak local_unnamed_addr addrspace(1) global [2 x %struct.StructTy1] [%struct.StructTy1 { i8* addrspacecast (i8 addrspace(4)* null to i8*), i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), i8 addrspace(2)* null, i8 addrspace(1)* null, i8 addrspace(4)* null }, %struct.StructTy1 { i8* addrspacecast (i8 addrspace(4)* null to i8*), i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), i8 addrspace(2)* null, i8 addrspace(1)* null, i8 addrspace(4)* null }], align 4 +StructTy1 A1[2]; + +// CHECK: @A2 = common local_unnamed_addr addrspace(1) global [2 x %struct.StructTy2] zeroinitializer, align 4 +StructTy2 A2[2]; + +// Test comparison with 0. + +// CHECK-LABEL: cmp_private +// CHECK: icmp eq i8* %p, addrspacecast (i8 addrspace(4)* null to i8*) +void cmp_private(private char* p) { + if (p != 0) + *p = 0; +} + +// CHECK-LABEL: cmp_local +// CHECK: icmp eq i8 addrspace(3)* %p, addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*) +void cmp_local(local char* p) { + if (p != 0) + *p = 0; +} + +// CHECK-LABEL: cmp_global +// CHECK: icmp eq i8 addrspace(1)* %p, null +void cmp_global(global char* p) { + if (p != 0) + *p = 0; +} + +// CHECK-LABEL: cmp_constant +// CHECK: icmp eq i8 addrspace(2)* %p, null +char cmp_constant(constant char* p) { + if (p != 0) + return *p; + else + return 0; +} + +// CHECK-LABEL: cmp_generic +// CHECK: icmp eq i8 addrspace(4)* %p, null +void cmp_generic(generic char* p) { + if (p != 0) + *p = 0; +} + +// Test comparison with NULL. + +// CHECK-LABEL: cmp_NULL_private +// CHECK: icmp eq i8* %p, addrspacecast (i8 addrspace(4)* null to i8*) +void cmp_NULL_private(private char* p) { + if (p != NULL) + *p = 0; +} + +// CHECK-LABEL: cmp_NULL_local +// CHECK: icmp eq i8 addrspace(3)* %p, addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*) +void cmp_NULL_local(local char* p) { + if (p != NULL) + *p = 0; +} + +// CHECK-LABEL: cmp_NULL_global +// CHECK: icmp eq i8 addrspace(1)* %p, null +void cmp_NULL_global(global char* p) { + if (p != NULL) + *p = 0; +} + +// CHECK-LABEL: cmp_NULL_constant +// CHECK: icmp eq i8 addrspace(2)* %p, null +char cmp_NULL_constant(constant char* p) { + if (p != NULL) + return *p; + else + return 0; +} + +// CHECK-LABEL: cmp_NULL_generic +// CHECK: icmp eq i8 addrspace(4)* %p, null +void cmp_NULL_generic(generic char* p) { + if (p != NULL) + *p = 0; +} + +// Test storage 0 as null pointer. +// CHECK-LABEL: test_storage_null_pointer +// CHECK: store i8* addrspacecast (i8 addrspace(4)* null to i8*), i8* addrspace(4)* %arg_private +// CHECK: store i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), i8 addrspace(3)* addrspace(4)* %arg_local +// CHECK: store i8 addrspace(1)* null, i8 addrspace(1)* addrspace(4)* %arg_global +// CHECK: store i8 addrspace(2)* null, i8 addrspace(2)* addrspace(4)* %arg_constant +// CHECK: store i8 addrspace(4)* null, i8 addrspace(4)* addrspace(4)* %arg_generic +void test_storage_null_pointer(private char** arg_private, + local char** arg_local, + global char** arg_global, + constant char** arg_constant, + generic char** arg_generic) { + *arg_private = 0; + *arg_local = 0; + *arg_global = 0; + *arg_constant = 0; + *arg_generic = 0; +} + +// Test storage NULL as null pointer. +// CHECK-LABEL: test_storage_null_pointer_NULL +// CHECK: store i8* addrspacecast (i8 addrspace(4)* null to i8*), i8* addrspace(4)* %arg_private +// CHECK: store i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), i8 addrspace(3)* addrspace(4)* %arg_local +// CHECK: store i8 addrspace(1)* null, i8 addrspace(1)* addrspace(4)* %arg_global +// CHECK: store i8 addrspace(2)* null, i8 addrspace(2)* addrspace(4)* %arg_constant +// CHECK: store i8 addrspace(4)* null, i8 addrspace(4)* addrspace(4)* %arg_generic +void test_storage_null_pointer_NULL(private char** arg_private, + local char** arg_local, + global char** arg_global, + constant char** arg_constant, + generic char** arg_generic) { + *arg_private = NULL; + *arg_local = NULL; + *arg_global = NULL; + *arg_constant = NULL; + *arg_generic = NULL; +} + +// Test pass null pointer to function as argument. +void test_pass_null_pointer_arg_calee(private char* arg_private, + local char* arg_local, + global char* arg_global, + constant char* arg_constant, + generic char* arg_generic); + +// CHECK-LABEL: test_pass_null_pointer_arg +// CHECK: call void @test_pass_null_pointer_arg_calee(i8* addrspacecast (i8 addrspace(4)* null to i8*), i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), i8 addrspace(1)* null, i8 addrspace(2)* null, i8 addrspace(4)* null) +// CHECK: call void @test_pass_null_pointer_arg_calee(i8* addrspacecast (i8 addrspace(4)* null to i8*), i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), i8 addrspace(1)* null, i8 addrspace(2)* null, i8 addrspace(4)* null) +void test_pass_null_pointer_arg(void) { + test_pass_null_pointer_arg_calee(0, 0, 0, 0, 0); + test_pass_null_pointer_arg_calee(NULL, NULL, NULL, NULL, NULL); +} + +// Test cast null pointer to size_t. +void test_cast_null_pointer_to_sizet_calee(size_t arg_private, + size_t arg_local, + size_t arg_global, + size_t arg_constant, + size_t arg_generic); + +// CHECK-LABEL: test_cast_null_pointer_to_sizet +// CHECK: call void @test_cast_null_pointer_to_sizet_calee(i64 ptrtoint (i8* addrspacecast (i8 addrspace(4)* null to i8*) to i64), i64 ptrtoint (i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*) to i64), i64 0, i64 0, i64 0) +// CHECK: call void @test_cast_null_pointer_to_sizet_calee(i64 ptrtoint (i8* addrspacecast (i8 addrspace(4)* null to i8*) to i64), i64 ptrtoint (i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*) to i64), i64 0, i64 0, i64 0) +void test_cast_null_pointer_to_sizet(void) { + test_cast_null_pointer_to_sizet_calee((size_t)((private char*)0), + (size_t)((local char*)0), + (size_t)((global char*)0), + (size_t)((constant char*)0), + (size_t)((generic char*)0)); + test_cast_null_pointer_to_sizet_calee((size_t)((private char*)NULL), + (size_t)((local char*)NULL), + (size_t)((global char*)NULL), + (size_t)((constant char*)0), // NULL cannot be casted to constant pointer since it is defined as a generic pointer + (size_t)((generic char*)NULL)); +} + +// Test comparision between null pointers. +#define TEST_EQ00(addr1, addr2) int test_eq00_##addr1##_##addr2(void) { return (addr1 char*)0 == (addr2 char*)0; } +#define TEST_EQ0N(addr1, addr2) int test_eq0N_##addr1##_##addr2(void) { return (addr1 char*)0 == (addr2 char*)NULL; } +#define TEST_EQN0(addr1, addr2) int test_eqN0_##addr1##_##addr2(void) { return (addr1 char*)NULL == (addr2 char*)0; } +#define TEST_EQNN(addr1, addr2) int test_eqNN_##addr1##_##addr2(void) { return (addr1 char*)0 == (addr2 char*)NULL; } +#define TEST_NE00(addr1, addr2) int test_ne00_##addr1##_##addr2(void) { return (addr1 char*)0 != (addr2 char*)0; } +#define TEST_NE0N(addr1, addr2) int test_ne0N_##addr1##_##addr2(void) { return (addr1 char*)0 != (addr2 char*)NULL; } +#define TEST_NEN0(addr1, addr2) int test_neN0_##addr1##_##addr2(void) { return (addr1 char*)NULL != (addr2 char*)0; } +#define TEST_NENN(addr1, addr2) int test_neNN_##addr1##_##addr2(void) { return (addr1 char*)0 != (addr2 char*)NULL; } +#define TEST(addr1, addr2) \ + TEST_EQ00(addr1, addr2) \ + TEST_EQ0N(addr1, addr2) \ + TEST_EQN0(addr1, addr2) \ + TEST_EQNN(addr1, addr2) \ + TEST_NE00(addr1, addr2) \ + TEST_NE0N(addr1, addr2) \ + TEST_NEN0(addr1, addr2) \ + TEST_NENN(addr1, addr2) + +// CHECK-LABEL: test_eq00_generic_private +// CHECK: ret i32 1 +// CHECK-LABEL: test_eq0N_generic_private +// CHECK: ret i32 1 +// CHECK-LABEL: test_eqN0_generic_private +// CHECK: ret i32 1 +// CHECK-LABEL: test_eqNN_generic_private +// CHECK: ret i32 1 +// CHECK-LABEL: test_ne00_generic_private +// CHECK: ret i32 0 +// CHECK-LABEL: test_ne0N_generic_private +// CHECK: ret i32 0 +// CHECK-LABEL: test_neN0_generic_private +// CHECK: ret i32 0 +// CHECK-LABEL: test_neNN_generic_private +// CHECK: ret i32 0 +TEST(generic, private) + +// CHECK-LABEL: test_eq00_generic_local +// CHECK: ret i32 1 +// CHECK-LABEL: test_eq0N_generic_local +// CHECK: ret i32 1 +// CHECK-LABEL: test_eqN0_generic_local +// CHECK: ret i32 1 +// CHECK-LABEL: test_eqNN_generic_local +// CHECK: ret i32 1 +// CHECK-LABEL: test_ne00_generic_local +// CHECK: ret i32 0 +// CHECK-LABEL: test_ne0N_generic_local +// CHECK: ret i32 0 +// CHECK-LABEL: test_neN0_generic_local +// CHECK: ret i32 0 +// CHECK-LABEL: test_neNN_generic_local +// CHECK: ret i32 0 +TEST(generic, local) + +// CHECK-LABEL: test_eq00_generic_global +// CHECK: ret i32 1 +// CHECK-LABEL: test_eq0N_generic_global +// CHECK: ret i32 1 +// CHECK-LABEL: test_eqN0_generic_global +// CHECK: ret i32 1 +// CHECK-LABEL: test_eqNN_generic_global +// CHECK: ret i32 1 +// CHECK-LABEL: test_ne00_generic_global +// CHECK: ret i32 0 +// CHECK-LABEL: test_ne0N_generic_global +// CHECK: ret i32 0 +// CHECK-LABEL: test_neN0_generic_global +// CHECK: ret i32 0 +// CHECK-LABEL: test_neNN_generic_global +// CHECK: ret i32 0 +TEST(generic, global) + +// CHECK-LABEL: test_eq00_generic_generic +// CHECK: ret i32 1 +// CHECK-LABEL: test_eq0N_generic_generic +// CHECK: ret i32 1 +// CHECK-LABEL: test_eqN0_generic_generic +// CHECK: ret i32 1 +// CHECK-LABEL: test_eqNN_generic_generic +// CHECK: ret i32 1 +// CHECK-LABEL: test_ne00_generic_generic +// CHECK: ret i32 0 +// CHECK-LABEL: test_ne0N_generic_generic +// CHECK: ret i32 0 +// CHECK-LABEL: test_neN0_generic_generic +// CHECK: ret i32 0 +// CHECK-LABEL: test_neNN_generic_generic +// CHECK: ret i32 0 +TEST(generic, generic) + +// CHECK-LABEL: test_eq00_constant_constant +// CHECK: ret i32 1 +TEST_EQ00(constant, constant) + +// Test cast to bool. + +// CHECK-LABEL: cast_bool_private +// CHECK: icmp eq i8* %p, addrspacecast (i8 addrspace(4)* null to i8*) +void cast_bool_private(private char* p) { + if (p) + *p = 0; +} + +// CHECK-LABEL: cast_bool_local +// CHECK: icmp eq i8 addrspace(3)* %p, addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*) +void cast_bool_local(local char* p) { + if (p) + *p = 0; +} + +// CHECK-LABEL: cast_bool_global +// CHECK: icmp eq i8 addrspace(1)* %p, null +void cast_bool_global(global char* p) { + if (p) + *p = 0; +} + +// CHECK-LABEL: cast_bool_constant +// CHECK: icmp eq i8 addrspace(2)* %p, null +char cast_bool_constant(constant char* p) { + if (p) + return *p; + else + return 0; +} + +// CHECK-LABEL: cast_bool_generic +// CHECK: icmp eq i8 addrspace(4)* %p, null +void cast_bool_generic(generic char* p) { + if (p) + *p = 0; +} + +// Test initialize a struct using memset. +// For large structures which is mostly zero, clang generats llvm.memset for +// the zero part and store for non-zero members. +typedef struct { + long a, b, c, d; + private char *p; +} StructTy3; + +// CHECK-LABEL: test_memset +// CHECK: call void @llvm.memset.p0i8.i64(i8* {{.*}}, i8 0, i64 32, i32 8, i1 false) +// CHECK: store i8* addrspacecast (i8 addrspace(4)* null to i8*), i8** {{.*}} +StructTy3 test_memset(void) { + StructTy3 S3 = {0, 0, 0, 0, 0}; + return S3; +} + +// Test casting literal 0 to pointer. +// A 0 literal casted to pointer should become a null pointer. + +// CHECK-LABEL: test_cast_0_to_ptr +// CHECK: ret i32* addrspacecast (i32 addrspace(4)* null to i32*) +private int* test_cast_0_to_ptr(void) { + return (private int*)0; +} + +// Test casting non-literal integer with 0 value to pointer. +// A non-literal integer expression with 0 value is casted to a pointer with +// zero value. + +// CHECK-LABEL: test_cast_int_to_ptr1 +// CHECK: ret i32* null +private int* test_cast_int_to_ptr1(void) { + return (private int*)((void)0, 0); +} + +// CHECK-LABEL: test_cast_int_to_ptr2 +// CHECK: ret i32* null +private int* test_cast_int_to_ptr2(void) { + int x = 0; + return (private int*)x; +} + +// Test logical operations. +// CHECK-LABEL: test_not_nullptr +// CHECK: ret i32 1 +int test_not_nullptr(void) { + return !(private char*)NULL; +} + +// CHECK-LABEL: test_and_nullptr +// CHECK: ret i32 0 +int test_and_nullptr(int a) { + return a && ((private char*)NULL); +} + +// CHECK-LABEL: test_not_ptr +// CHECK: %[[lnot:.*]] = icmp eq i8* %p, addrspacecast (i8 addrspace(4)* null to i8*) +// CHECK: %[[lnot_ext:.*]] = zext i1 %[[lnot]] to i32 +// CHECK: ret i32 %[[lnot_ext]] +int test_not_ptr(private char* p) { + return !p; +} +// CHECK-LABEL: test_and_ptr +// CHECK: %[[tobool:.*]] = icmp ne i8* %p1, addrspacecast (i8 addrspace(4)* null to i8*) +// CHECK: %[[tobool1:.*]] = icmp ne i8 addrspace(3)* %p2, addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*) +// CHECK: %[[res:.*]] = and i1 %[[tobool]], %[[tobool1]] +// CHECK: %[[land_ext:.*]] = zext i1 %[[res]] to i32 +// CHECK: ret i32 %[[land_ext]] +int test_and_ptr(private char* p1, local char* p2) { + return p1 && p2; +} + +// Test folding of null pointer in function scope. +// NOOPT-LABEL: test_fold +// NOOPT: call void @test_fold_callee +// NOOPT: store i32 addrspace(1)* null, i32 addrspace(1)** %glob, align 4 +// NOOPT: %{{.*}} = sub i64 %{{.*}}, 0 +// NOOPT: call void @test_fold_callee +// NOOPT: %{{.*}} = add nsw i64 %{{.*}}, sext (i32 ptrtoint (i32* addrspacecast (i32 addrspace(4)* null to i32*) to i32) to i64) +// NOOPT: %{{.*}} = sub nsw i64 %{{.*}}, 1 +void test_fold_callee(void); +void test_fold(void) { + global int* glob = (test_fold_callee(), (global int*)(generic char*)0); + long x = glob - (global int*)(generic char*)0; + x = x + (int)(test_fold_callee(), (private int*)(generic char*)(global short*)0); + x = x - (int)((private int*)0 == (private int*)(generic char*)0); +} diff --git a/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl b/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl deleted file mode 100644 index 589d00d1eaa49..0000000000000 --- a/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl +++ /dev/null @@ -1,48 +0,0 @@ -// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -O0 -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -verify -o - %s | FileCheck -check-prefix=X86 %s - -// Make sure this is silently accepted on other targets. - -__attribute__((amdgpu_num_vgpr(64))) // expected-no-diagnostics -kernel void test_num_vgpr64() { -// CHECK: define amdgpu_kernel void @test_num_vgpr64() [[ATTR_VGPR64:#[0-9]+]] -} - -__attribute__((amdgpu_num_sgpr(32))) // expected-no-diagnostics -kernel void test_num_sgpr32() { -// CHECK: define amdgpu_kernel void @test_num_sgpr32() [[ATTR_SGPR32:#[0-9]+]] -} - -__attribute__((amdgpu_num_vgpr(64), amdgpu_num_sgpr(32))) // expected-no-diagnostics -kernel void test_num_vgpr64_sgpr32() { -// CHECK: define amdgpu_kernel void @test_num_vgpr64_sgpr32() [[ATTR_VGPR64_SGPR32:#[0-9]+]] - -} - -__attribute__((amdgpu_num_sgpr(20), amdgpu_num_vgpr(40))) // expected-no-diagnostics -kernel void test_num_sgpr20_vgpr40() { -// CHECK: define amdgpu_kernel void @test_num_sgpr20_vgpr40() [[ATTR_SGPR20_VGPR40:#[0-9]+]] -} - -__attribute__((amdgpu_num_vgpr(0))) // expected-no-diagnostics -kernel void test_num_vgpr0() { -} - -__attribute__((amdgpu_num_sgpr(0))) // expected-no-diagnostics -kernel void test_num_sgpr0() { -} - -__attribute__((amdgpu_num_vgpr(0), amdgpu_num_sgpr(0))) // expected-no-diagnostics -kernel void test_num_vgpr0_sgpr0() { -} - - -// X86-NOT: "amdgpu_num_vgpr" -// X86-NOT: "amdgpu_num_sgpr" - -// CHECK-NOT: "amdgpu_num_vgpr"="0" -// CHECK-NOT: "amdgpu_num_sgpr"="0" -// CHECK-DAG: attributes [[ATTR_VGPR64]] = { nounwind "amdgpu_num_vgpr"="64" -// CHECK-DAG: attributes [[ATTR_SGPR32]] = { nounwind "amdgpu_num_sgpr"="32" -// CHECK-DAG: attributes [[ATTR_VGPR64_SGPR32]] = { nounwind "amdgpu_num_sgpr"="32" "amdgpu_num_vgpr"="64" -// CHECK-DAG: attributes [[ATTR_SGPR20_VGPR40]] = { nounwind "amdgpu_num_sgpr"="20" "amdgpu_num_vgpr"="40" diff --git a/test/CodeGenOpenCL/as_type.cl b/test/CodeGenOpenCL/as_type.cl index 7fc3b02bdc757..51e82e3fef8ed 100644 --- a/test/CodeGenOpenCL/as_type.cl +++ b/test/CodeGenOpenCL/as_type.cl @@ -51,7 +51,7 @@ int f6(char4 x) { return __builtin_astype(x, int); } -//CHECK: define spir_func <3 x i8> @f7(<3 x i8> %[[x:.*]]) +//CHECK: define spir_func <3 x i8> @f7(<3 x i8> returned %[[x:.*]]) //CHECK-NOT: bitcast //CHECK-NOT: shufflevector //CHECK: ret <3 x i8> %[[x]] @@ -66,3 +66,42 @@ char3 f7(char3 x) { int3 f8(char16 x) { return __builtin_astype(x, int3); } + +//CHECK: define spir_func i32 addrspace(1)* @addr_cast(i32* readnone %[[x:.*]]) +//CHECK: %[[cast:.*]] = addrspacecast i32* %[[x]] to i32 addrspace(1)* +//CHECK: ret i32 addrspace(1)* %[[cast]] +global int* addr_cast(int *x) { + return __builtin_astype(x, global int*); +} + +//CHECK: define spir_func i32 addrspace(1)* @int_to_ptr(i32 %[[x:.*]]) +//CHECK: %[[cast:.*]] = inttoptr i32 %[[x]] to i32 addrspace(1)* +//CHECK: ret i32 addrspace(1)* %[[cast]] +global int* int_to_ptr(int x) { + return __builtin_astype(x, global int*); +} + +//CHECK: define spir_func i32 @ptr_to_int(i32* %[[x:.*]]) +//CHECK: %[[cast:.*]] = ptrtoint i32* %[[x]] to i32 +//CHECK: ret i32 %[[cast]] +int ptr_to_int(int *x) { + return __builtin_astype(x, int); +} + +//CHECK: define spir_func <3 x i8> @ptr_to_char3(i32* %[[x:.*]]) +//CHECK: %[[cast1:.*]] = ptrtoint i32* %[[x]] to i32 +//CHECK: %[[cast2:.*]] = bitcast i32 %[[cast1]] to <4 x i8> +//CHECK: %[[astype:.*]] = shufflevector <4 x i8> %[[cast2]], <4 x i8> undef, <3 x i32> <i32 0, i32 1, i32 2> +//CHECK: ret <3 x i8> %[[astype]] +char3 ptr_to_char3(int *x) { + return __builtin_astype(x, char3); +} + +//CHECK: define spir_func i32* @char3_to_ptr(<3 x i8> %[[x:.*]]) +//CHECK: %[[astype:.*]] = shufflevector <3 x i8> %[[x]], <3 x i8> undef, <4 x i32> <i32 0, i32 1, i32 2, i32 undef> +//CHECK: %[[cast1:.*]] = bitcast <4 x i8> %[[astype]] to i32 +//CHECK: %[[cast2:.*]] = inttoptr i32 %[[cast1]] to i32* +//CHECK: ret i32* %[[cast2]] +int* char3_to_ptr(char3 x) { + return __builtin_astype(x, int*); +} diff --git a/test/CodeGenOpenCL/builtins-amdgcn-error.cl b/test/CodeGenOpenCL/builtins-amdgcn-error.cl deleted file mode 100644 index 89c3e490ecdd8..0000000000000 --- a/test/CodeGenOpenCL/builtins-amdgcn-error.cl +++ /dev/null @@ -1,18 +0,0 @@ -// REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-unknown-amdhsa -target-cpu tahiti -verify -S -o - %s - -// FIXME: We only get one error if the functions are the other order in the -// file. - -typedef unsigned long ulong; - -ulong test_s_memrealtime() -{ - return __builtin_amdgcn_s_memrealtime(); // expected-error {{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}} -} - -void test_s_sleep(int x) -{ - __builtin_amdgcn_s_sleep(x); // expected-error {{argument to '__builtin_amdgcn_s_sleep' must be a constant integer}} -} - diff --git a/test/CodeGenOpenCL/builtins-amdgcn-vi.cl b/test/CodeGenOpenCL/builtins-amdgcn-vi.cl index cda87a8e1ed17..0c88fdfb724d5 100644 --- a/test/CodeGenOpenCL/builtins-amdgcn-vi.cl +++ b/test/CodeGenOpenCL/builtins-amdgcn-vi.cl @@ -1,8 +1,79 @@ // REQUIRES: amdgpu-registered-target // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu tonga -S -emit-llvm -o - %s | FileCheck %s +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + typedef unsigned long ulong; +// CHECK-LABEL: @test_div_fixup_f16 +// CHECK: call half @llvm.amdgcn.div.fixup.f16 +void test_div_fixup_f16(global half* out, half a, half b, half c) +{ + *out = __builtin_amdgcn_div_fixuph(a, b, c); +} + +// CHECK-LABEL: @test_rcp_f16 +// CHECK: call half @llvm.amdgcn.rcp.f16 +void test_rcp_f16(global half* out, half a) +{ + *out = __builtin_amdgcn_rcph(a); +} + +// CHECK-LABEL: @test_rsq_f16 +// CHECK: call half @llvm.amdgcn.rsq.f16 +void test_rsq_f16(global half* out, half a) +{ + *out = __builtin_amdgcn_rsqh(a); +} + +// CHECK-LABEL: @test_sin_f16 +// CHECK: call half @llvm.amdgcn.sin.f16 +void test_sin_f16(global half* out, half a) +{ + *out = __builtin_amdgcn_sinh(a); +} + +// CHECK-LABEL: @test_cos_f16 +// CHECK: call half @llvm.amdgcn.cos.f16 +void test_cos_f16(global half* out, half a) +{ + *out = __builtin_amdgcn_cosh(a); +} + +// CHECK-LABEL: @test_ldexp_f16 +// CHECK: call half @llvm.amdgcn.ldexp.f16 +void test_ldexp_f16(global half* out, half a, int b) +{ + *out = __builtin_amdgcn_ldexph(a, b); +} + +// CHECK-LABEL: @test_frexp_mant_f16 +// CHECK: call half @llvm.amdgcn.frexp.mant.f16 +void test_frexp_mant_f16(global half* out, half a) +{ + *out = __builtin_amdgcn_frexp_manth(a); +} + +// CHECK-LABEL: @test_frexp_exp_f16 +// CHECK: call i16 @llvm.amdgcn.frexp.exp.i16.f16 +void test_frexp_exp_f16(global short* out, half a) +{ + *out = __builtin_amdgcn_frexp_exph(a); +} + +// CHECK-LABEL: @test_fract_f16 +// CHECK: call half @llvm.amdgcn.fract.f16 +void test_fract_f16(global half* out, half a) +{ + *out = __builtin_amdgcn_fracth(a); +} + +// CHECK-LABEL: @test_class_f16 +// CHECK: call i1 @llvm.amdgcn.class.f16 +void test_class_f16(global half* out, half a, int b) +{ + *out = __builtin_amdgcn_classh(a, b); +} // CHECK-LABEL: @test_s_memrealtime // CHECK: call i64 @llvm.amdgcn.s.memrealtime() diff --git a/test/CodeGenOpenCL/builtins-amdgcn.cl b/test/CodeGenOpenCL/builtins-amdgcn.cl index 6cac2a44bc79a..ce566b6e876a1 100644 --- a/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -4,6 +4,7 @@ #pragma OPENCL EXTENSION cl_khr_fp64 : enable typedef unsigned long ulong; +typedef unsigned int uint; // CHECK-LABEL: @test_div_scale_f64 // CHECK: call { double, i1 } @llvm.amdgcn.div.scale.f64(double %a, double %b, i1 true) @@ -165,14 +166,14 @@ void test_frexp_mant_f64(global double* out, double a) } // CHECK-LABEL: @test_frexp_exp_f32 -// CHECK: call i32 @llvm.amdgcn.frexp.exp.f32 +// CHECK: call i32 @llvm.amdgcn.frexp.exp.i32.f32 void test_frexp_exp_f32(global int* out, float a) { *out = __builtin_amdgcn_frexp_expf(a); } // CHECK-LABEL: @test_frexp_exp_f64 -// CHECK: call i32 @llvm.amdgcn.frexp.exp.f64 +// CHECK: call i32 @llvm.amdgcn.frexp.exp.i32.f64 void test_frexp_exp_f64(global int* out, double a) { *out = __builtin_amdgcn_frexp_exp(a); @@ -199,6 +200,55 @@ void test_lerp(global int* out, int a, int b, int c) *out = __builtin_amdgcn_lerp(a, b, c); } +// CHECK-LABEL: @test_sicmp_i32 +// CHECK: call i64 @llvm.amdgcn.icmp.i32(i32 %a, i32 %b, i32 32) +void test_sicmp_i32(global ulong* out, int a, int b) +{ + *out = __builtin_amdgcn_sicmp(a, b, 32); +} + +// CHECK-LABEL: @test_uicmp_i32 +// CHECK: call i64 @llvm.amdgcn.icmp.i32(i32 %a, i32 %b, i32 32) +void test_uicmp_i32(global ulong* out, uint a, uint b) +{ + *out = __builtin_amdgcn_uicmp(a, b, 32); +} + +// CHECK-LABEL: @test_sicmp_i64 +// CHECK: call i64 @llvm.amdgcn.icmp.i64(i64 %a, i64 %b, i32 38) +void test_sicmp_i64(global ulong* out, long a, long b) +{ + *out = __builtin_amdgcn_sicmpl(a, b, 39-1); +} + +// CHECK-LABEL: @test_uicmp_i64 +// CHECK: call i64 @llvm.amdgcn.icmp.i64(i64 %a, i64 %b, i32 35) +void test_uicmp_i64(global ulong* out, ulong a, ulong b) +{ + *out = __builtin_amdgcn_uicmpl(a, b, 30+5); +} + +// CHECK-LABEL: @test_ds_swizzle +// CHECK: call i32 @llvm.amdgcn.ds.swizzle(i32 %a, i32 32) +void test_ds_swizzle(global int* out, int a) +{ + *out = __builtin_amdgcn_ds_swizzle(a, 32); +} + +// CHECK-LABEL: @test_fcmp_f32 +// CHECK: call i64 @llvm.amdgcn.fcmp.f32(float %a, float %b, i32 5) +void test_fcmp_f32(global ulong* out, float a, float b) +{ + *out = __builtin_amdgcn_fcmpf(a, b, 5); +} + +// CHECK-LABEL: @test_fcmp_f64 +// CHECK: call i64 @llvm.amdgcn.fcmp.f64(double %a, double %b, i32 6) +void test_fcmp_f64(global ulong* out, double a, double b) +{ + *out = __builtin_amdgcn_fcmp(a, b, 3+3); +} + // CHECK-LABEL: @test_class_f32 // CHECK: call i1 @llvm.amdgcn.class.f32 void test_class_f32(global float* out, float a, int b) @@ -220,6 +270,13 @@ void test_s_barrier() __builtin_amdgcn_s_barrier(); } +// CHECK-LABEL: @test_wave_barrier +// CHECK: call void @llvm.amdgcn.wave.barrier( +void test_wave_barrier() +{ + __builtin_amdgcn_wave_barrier(); +} + // CHECK-LABEL: @test_s_memtime // CHECK: call i64 @llvm.amdgcn.s.memtime() void test_s_memtime(global ulong* out) @@ -236,6 +293,24 @@ void test_s_sleep() __builtin_amdgcn_s_sleep(15); } +// CHECK-LABEL: @test_s_incperflevel +// CHECK: call void @llvm.amdgcn.s.incperflevel(i32 1) +// CHECK: call void @llvm.amdgcn.s.incperflevel(i32 15) +void test_s_incperflevel() +{ + __builtin_amdgcn_s_incperflevel(1); + __builtin_amdgcn_s_incperflevel(15); +} + +// CHECK-LABEL: @test_s_decperflevel +// CHECK: call void @llvm.amdgcn.s.decperflevel(i32 1) +// CHECK: call void @llvm.amdgcn.s.decperflevel(i32 15) +void test_s_decperflevel() +{ + __builtin_amdgcn_s_decperflevel(1); + __builtin_amdgcn_s_decperflevel(15); +} + // CHECK-LABEL: @test_cubeid( // CHECK: call float @llvm.amdgcn.cubeid(float %a, float %b, float %c) void test_cubeid(global float* out, float a, float b, float c) { diff --git a/test/CodeGenOpenCL/cast_image.cl b/test/CodeGenOpenCL/cast_image.cl new file mode 100644 index 0000000000000..479404a91371d --- /dev/null +++ b/test/CodeGenOpenCL/cast_image.cl @@ -0,0 +1,18 @@ +// RUN: %clang_cc1 -emit-llvm -o - -triple amdgcn--amdhsa %s | FileCheck --check-prefix=AMDGCN %s +// RUN: %clang_cc1 -emit-llvm -o - -triple spir-unknown-unknown %s | FileCheck --check-prefix=SPIR %s + +#ifdef __AMDGCN__ + +constant int* convert(image2d_t img) { + // AMDGCN: bitcast %opencl.image2d_ro_t addrspace(2)* %img to i32 addrspace(2)* + return __builtin_astype(img, constant int*); +} + +#else + +global int* convert(image2d_t img) { + // SPIR: bitcast %opencl.image2d_ro_t addrspace(1)* %img to i32 addrspace(1)* + return __builtin_astype(img, global int*); +} + +#endif diff --git a/test/CodeGenOpenCL/cl20-device-side-enqueue.cl b/test/CodeGenOpenCL/cl20-device-side-enqueue.cl index 783ce0254a1af..add13b9e3fc47 100644 --- a/test/CodeGenOpenCL/cl20-device-side-enqueue.cl +++ b/test/CodeGenOpenCL/cl20-device-side-enqueue.cl @@ -1,110 +1,147 @@ -// RUN: %clang_cc1 %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir-unknown-unknown" | FileCheck %s --check-prefix=COMMON --check-prefix=B32 +// RUN: %clang_cc1 %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefix=COMMON --check-prefix=B64 typedef void (^bl_t)(local void *); +// N.B. The check here only exists to set BL_GLOBAL +// COMMON: @block_G = {{.*}}bitcast ([[BL_GLOBAL:[^@]+@__block_literal_global(\.[0-9]+)?]] const bl_t block_G = (bl_t) ^ (local void *a) {}; kernel void device_side_enqueue(global int *a, global int *b, int i) { - // CHECK: %default_queue = alloca %opencl.queue_t* + // COMMON: %default_queue = alloca %opencl.queue_t* queue_t default_queue; - // CHECK: %flags = alloca i32 + // COMMON: %flags = alloca i32 unsigned flags = 0; - // CHECK: %ndrange = alloca %opencl.ndrange_t* + // COMMON: %ndrange = alloca %opencl.ndrange_t* ndrange_t ndrange; - // CHECK: %clk_event = alloca %opencl.clk_event_t* + // COMMON: %clk_event = alloca %opencl.clk_event_t* clk_event_t clk_event; - // CHECK: %event_wait_list = alloca %opencl.clk_event_t* + // COMMON: %event_wait_list = alloca %opencl.clk_event_t* clk_event_t event_wait_list; - // CHECK: %event_wait_list2 = alloca [1 x %opencl.clk_event_t*] + // COMMON: %event_wait_list2 = alloca [1 x %opencl.clk_event_t*] clk_event_t event_wait_list2[] = {clk_event}; - // CHECK: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t*, %opencl.queue_t** %default_queue - // CHECK: [[FLAGS:%[0-9]+]] = load i32, i32* %flags - // CHECK: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange - // CHECK: [[BL:%[0-9]+]] = bitcast <{ i8*, i32, i32, i8*, %struct.__block_descriptor addrspace(3)*, i32{{.*}}, i32{{.*}}, i32{{.*}} }>* %block to void ()* - // CHECK: [[BL_I8:%[0-9]+]] = bitcast void ()* [[BL]] to i8* - // CHECK: call i32 @__enqueue_kernel_basic(%opencl.queue_t* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* [[BL_I8]]) + // COMMON: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t{{.*}}*, %opencl.queue_t{{.*}}** %default_queue + // COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags + // COMMON: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange + // COMMON: [[BL:%[0-9]+]] = bitcast <{ i8*, i32, i32, i8*, %struct.__block_descriptor addrspace(2)*, i32{{.*}}, i32{{.*}}, i32{{.*}} }>* %block to void ()* + // COMMON: [[BL_I8:%[0-9]+]] = bitcast void ()* [[BL]] to i8* + // COMMON: call i32 @__enqueue_kernel_basic(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* [[BL_I8]]) enqueue_kernel(default_queue, flags, ndrange, ^(void) { a[i] = b[i]; }); - // CHECK: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t*, %opencl.queue_t** %default_queue - // CHECK: [[FLAGS:%[0-9]+]] = load i32, i32* %flags - // CHECK: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange - // CHECK: [[BL:%[0-9]+]] = bitcast <{ i8*, i32, i32, i8*, %struct.__block_descriptor addrspace(3)*, i32{{.*}}, i32{{.*}}, i32{{.*}} }>* %block3 to void ()* - // CHECK: [[BL_I8:%[0-9]+]] = bitcast void ()* [[BL]] to i8* - // CHECK: call i32 @__enqueue_kernel_basic_events(%opencl.queue_t* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i32 2, %opencl.clk_event_t** %event_wait_list, %opencl.clk_event_t** %clk_event, i8* [[BL_I8]]) + // COMMON: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t{{.*}}*, %opencl.queue_t{{.*}}** %default_queue + // COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags + // COMMON: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange + // COMMON: [[WAIT_EVNT:%[0-9]+]] = addrspacecast %opencl.clk_event_t{{.*}}** %event_wait_list to %opencl.clk_event_t{{.*}}* addrspace(4)* + // COMMON: [[EVNT:%[0-9]+]] = addrspacecast %opencl.clk_event_t{{.*}}** %clk_event to %opencl.clk_event_t{{.*}}* addrspace(4)* + // COMMON: [[BL:%[0-9]+]] = bitcast <{ i8*, i32, i32, i8*, %struct.__block_descriptor addrspace(2)*, i32{{.*}}, i32{{.*}}, i32{{.*}} }>* %block3 to void ()* + // COMMON: [[BL_I8:%[0-9]+]] = bitcast void ()* [[BL]] to i8* + // COMMON: call i32 @__enqueue_kernel_basic_events(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* {{.*}}, i32 2, %opencl.clk_event_t{{.*}}* addrspace(4)* [[WAIT_EVNT]], %opencl.clk_event_t{{.*}}* addrspace(4)* [[EVNT]], i8* [[BL_I8]]) enqueue_kernel(default_queue, flags, ndrange, 2, &event_wait_list, &clk_event, ^(void) { a[i] = b[i]; }); - // CHECK: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t*, %opencl.queue_t** %default_queue - // CHECK: [[FLAGS:%[0-9]+]] = load i32, i32* %flags - // CHECK: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange - // CHECK: call i32 (%opencl.queue_t*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(3)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 256) + // COMMON: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t{{.*}}*, %opencl.queue_t{{.*}}** %default_queue + // COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags + // COMMON: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange + // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 256) + // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i64 256) enqueue_kernel(default_queue, flags, ndrange, ^(local void *p) { return; }, 256); char c; - // CHECK: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t*, %opencl.queue_t** %default_queue - // CHECK: [[FLAGS:%[0-9]+]] = load i32, i32* %flags - // CHECK: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange - // CHECK: [[SIZE:%[0-9]+]] = zext i8 {{%[0-9]+}} to i32 - // CHECK: call i32 (%opencl.queue_t*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(3)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 [[SIZE]]) + // COMMON: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t{{.*}}*, %opencl.queue_t{{.*}}** %default_queue + // COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags + // COMMON: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange + // B32: [[SIZE:%[0-9]+]] = zext i8 {{%[0-9]+}} to i32 + // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 [[SIZE]]) + // B64: [[SIZE:%[0-9]+]] = zext i8 {{%[0-9]+}} to i64 + // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i64 [[SIZE]]) enqueue_kernel(default_queue, flags, ndrange, ^(local void *p) { return; }, c); - // CHECK: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t*, %opencl.queue_t** %default_queue - // CHECK: [[FLAGS:%[0-9]+]] = load i32, i32* %flags - // CHECK: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange - // CHECK: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x %opencl.clk_event_t*], [1 x %opencl.clk_event_t*]* %event_wait_list2, i32 0, i32 0 - // CHECK: call i32 (%opencl.queue_t*, i32, %opencl.ndrange_t*, i32, %opencl.clk_event_t**, %opencl.clk_event_t**, i8*, i32, ...) @__enqueue_kernel_events_vaargs(%opencl.queue_t* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i32 2, %opencl.clk_event_t** [[AD]], %opencl.clk_event_t** %clk_event, i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(3)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 256) + // COMMON: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t{{.*}}*, %opencl.queue_t{{.*}}** %default_queue + // COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags + // COMMON: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange + // COMMON: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x %opencl.clk_event_t*], [1 x %opencl.clk_event_t*]* %event_wait_list2, i32 0, i32 0 + // COMMON: [[WAIT_EVNT:%[0-9]+]] = addrspacecast %opencl.clk_event_t{{.*}}** [[AD]] to %opencl.clk_event_t{{.*}}* addrspace(4)* + // COMMON: [[EVNT:%[0-9]+]] = addrspacecast %opencl.clk_event_t{{.*}}** %clk_event to %opencl.clk_event_t{{.*}}* addrspace(4)* + // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i32, %opencl.clk_event_t{{.*}}* addrspace(4)*, %opencl.clk_event_t{{.*}}* addrspace(4)*, i8*, i32, ...) @__enqueue_kernel_events_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* {{.*}}, i32 2, %opencl.clk_event_t{{.*}} [[WAIT_EVNT]], %opencl.clk_event_t{{.*}} [[EVNT]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 256) + // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i32, %opencl.clk_event_t{{.*}}* addrspace(4)*, %opencl.clk_event_t{{.*}}* addrspace(4)*, i8*, i32, ...) @__enqueue_kernel_events_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* {{.*}}, i32 2, %opencl.clk_event_t{{.*}} [[WAIT_EVNT]], %opencl.clk_event_t{{.*}} [[EVNT]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i64 256) enqueue_kernel(default_queue, flags, ndrange, 2, event_wait_list2, &clk_event, ^(local void *p) { return; }, 256); - // CHECK: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t*, %opencl.queue_t** %default_queue - // CHECK: [[FLAGS:%[0-9]+]] = load i32, i32* %flags - // CHECK: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange - // CHECK: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x %opencl.clk_event_t*], [1 x %opencl.clk_event_t*]* %event_wait_list2, i32 0, i32 0 - // CHECK: [[SIZE:%[0-9]+]] = zext i8 {{%[0-9]+}} to i32 - // CHECK: call i32 (%opencl.queue_t*, i32, %opencl.ndrange_t*, i32, %opencl.clk_event_t**, %opencl.clk_event_t**, i8*, i32, ...) @__enqueue_kernel_events_vaargs(%opencl.queue_t* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i32 2, %opencl.clk_event_t** [[AD]], %opencl.clk_event_t** %clk_event, i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(3)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 [[SIZE]]) + // COMMON: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t{{.*}}*, %opencl.queue_t{{.*}}** %default_queue + // COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags + // COMMON: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange + // COMMON: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x %opencl.clk_event_t*], [1 x %opencl.clk_event_t*]* %event_wait_list2, i32 0, i32 0 + // COMMON: [[WAIT_EVNT:%[0-9]+]] = addrspacecast %opencl.clk_event_t{{.*}}** [[AD]] to %opencl.clk_event_t{{.*}}* addrspace(4)* + // COMMON: [[EVNT:%[0-9]+]] = addrspacecast %opencl.clk_event_t{{.*}}** %clk_event to %opencl.clk_event_t{{.*}}* addrspace(4)* + // B32: [[SIZE:%[0-9]+]] = zext i8 {{%[0-9]+}} to i32 + // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i32, %opencl.clk_event_t{{.*}}* addrspace(4)*, %opencl.clk_event_t{{.*}}* addrspace(4)*, i8*, i32, ...) @__enqueue_kernel_events_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* {{.*}}, i32 2, %opencl.clk_event_t{{.*}}* addrspace(4)* [[WAIT_EVNT]], %opencl.clk_event_t{{.*}}* addrspace(4)* [[EVNT]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 [[SIZE]]) + // B64: [[SIZE:%[0-9]+]] = zext i8 {{%[0-9]+}} to i64 + // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i32, %opencl.clk_event_t{{.*}}* addrspace(4)*, %opencl.clk_event_t{{.*}}* addrspace(4)*, i8*, i32, ...) @__enqueue_kernel_events_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* {{.*}}, i32 2, %opencl.clk_event_t{{.*}}* addrspace(4)* [[WAIT_EVNT]], %opencl.clk_event_t{{.*}}* addrspace(4)* [[EVNT]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i64 [[SIZE]]) enqueue_kernel(default_queue, flags, ndrange, 2, event_wait_list2, &clk_event, ^(local void *p) { return; }, c); + long l; + // COMMON: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t{{.*}}*, %opencl.queue_t{{.*}}** %default_queue + // COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags + // COMMON: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange + // B32: [[SIZE:%[0-9]+]] = trunc i64 {{%[0-9]+}} to i32 + // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 [[SIZE]]) + // B64: [[SIZE:%[0-9]+]] = load i64, i64* %l + // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i64 [[SIZE]]) + enqueue_kernel(default_queue, flags, ndrange, + ^(local void *p) { + return; + }, + l); + + // COMMON: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t*, %opencl.queue_t** %default_queue + // COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags + // COMMON: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange + // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 0) + // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i64 4294967296) + enqueue_kernel(default_queue, flags, ndrange, + ^(local void *p) { + return; + }, + 4294967296L); + + // The full type of these expressions are long (and repeated elsewhere), so we + // capture it as part of the regex for convenience and clarity. + // COMMON: store void ()* bitcast ([[BL_A:[^@]+@__block_literal_global.[0-9]+]] to void ()*), void ()** %block_A void (^const block_A)(void) = ^{ return; }; + + // COMMON: store void (i8 addrspace(3)*)* bitcast ([[BL_B:[^@]+@__block_literal_global.[0-9]+]] to void (i8 addrspace(3)*)*), void (i8 addrspace(3)*)** %block_B void (^const block_B)(local void *) = ^(local void *a) { return; }; - // CHECK: [[BL:%[0-9]+]] = load void ()*, void ()** %block_A - // CHECK: [[BL_I8:%[0-9]+]] = bitcast void ()* [[BL]] to i8* - // CHECK: call i32 @__get_kernel_work_group_size_impl(i8* [[BL_I8]]) + // COMMON: call i32 @__get_kernel_work_group_size_impl(i8* bitcast ([[BL_A]] to i8*)) unsigned size = get_kernel_work_group_size(block_A); - // CHECK: [[BL:%[0-9]+]] = load void (i8 addrspace(2)*)*, void (i8 addrspace(2)*)** %block_B - // CHECK: [[BL_I8:%[0-9]+]] = bitcast void (i8 addrspace(2)*)* [[BL]] to i8* - // CHECK: call i32 @__get_kernel_work_group_size_impl(i8* [[BL_I8]]) + // COMMON: call i32 @__get_kernel_work_group_size_impl(i8* bitcast ([[BL_B]] to i8*)) size = get_kernel_work_group_size(block_B); - // CHECK: [[BL:%[0-9]+]] = load void ()*, void ()** %block_A - // CHECK: [[BL_I8:%[0-9]+]] = bitcast void ()* [[BL]] to i8* - // CHECK: call i32 @__get_kernel_preferred_work_group_multiple_impl(i8* [[BL_I8]]) + // COMMON: call i32 @__get_kernel_preferred_work_group_multiple_impl(i8* bitcast ([[BL_A]] to i8*)) size = get_kernel_preferred_work_group_size_multiple(block_A); - // CHECK: [[BL:%[0-9]+]] = load void (i8 addrspace(2)*)*, void (i8 addrspace(2)*)* addrspace(1)* @block_G - // CHECK: [[BL_I8:%[0-9]+]] = bitcast void (i8 addrspace(2)*)* [[BL]] to i8* - // CHECK: call i32 @__get_kernel_preferred_work_group_multiple_impl(i8* [[BL_I8]]) + // COMMON: call i32 @__get_kernel_preferred_work_group_multiple_impl(i8* bitcast ([[BL_GLOBAL]] to i8*)) size = get_kernel_preferred_work_group_size_multiple(block_G); } diff --git a/test/CodeGenOpenCL/const-str-array-decay.cl b/test/CodeGenOpenCL/const-str-array-decay.cl index af8c35542f3b0..353aa3a9a083d 100644 --- a/test/CodeGenOpenCL/const-str-array-decay.cl +++ b/test/CodeGenOpenCL/const-str-array-decay.cl @@ -6,6 +6,6 @@ kernel void str_array_decy() { test_func("Test string literal"); } -// CHECK: i8 addrspace(3)* getelementptr inbounds ([20 x i8], [20 x i8] addrspace(3)* +// CHECK: i8 addrspace(2)* getelementptr inbounds ([20 x i8], [20 x i8] addrspace(2)* // CHECK-NOT: addrspacecast diff --git a/test/CodeGenOpenCL/constant-addr-space-globals.cl b/test/CodeGenOpenCL/constant-addr-space-globals.cl index f81a51458c428..4f0d1ea23e564 100644 --- a/test/CodeGenOpenCL/constant-addr-space-globals.cl +++ b/test/CodeGenOpenCL/constant-addr-space-globals.cl @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 %s -cl-opt-disable -ffake-address-space-map -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple "spir64-unknown-unknown" -cl-opt-disable -ffake-address-space-map -emit-llvm -o - | FileCheck %s // CHECK: @array = addrspace({{[0-9]+}}) constant __constant float array[2] = {0.0f, 1.0f}; @@ -6,3 +6,22 @@ __constant float array[2] = {0.0f, 1.0f}; kernel void test(global float *out) { *out = array[0]; } + +// Test that we don't use directly initializers for const aggregates +// but create a copy in the original address space (unless a variable itself is +// in the constant address space). + +void foo(constant const int *p1, const int *p2, const int *p3); +// CHECK: @k.arr1 = internal addrspace(2) constant [3 x i32] [i32 1, i32 2, i32 3] +// CHECK: @k.arr2 = private unnamed_addr addrspace(2) constant [3 x i32] [i32 4, i32 5, i32 6] +// CHECK: @k.arr3 = private unnamed_addr addrspace(2) constant [3 x i32] [i32 7, i32 8, i32 9] +kernel void k(void) { + // CHECK-NOT: %arr1 = alloca [3 x i32] + constant const int arr1[] = {1, 2, 3}; + // CHECK: %arr2 = alloca [3 x i32] + const int arr2[] = {4, 5, 6}; + // CHECK: %arr3 = alloca [3 x i32] + int arr3[] = {7, 8, 9}; + + foo(arr1, arr2, arr3); +} diff --git a/test/CodeGenOpenCL/convergent.cl b/test/CodeGenOpenCL/convergent.cl new file mode 100644 index 0000000000000..c6bcb52a6d750 --- /dev/null +++ b/test/CodeGenOpenCL/convergent.cl @@ -0,0 +1,118 @@ +// RUN: %clang_cc1 -triple spir-unknown-unknown -emit-llvm %s -o - | opt -instnamer -S | FileCheck %s + +void convfun(void) __attribute__((convergent)); +void non_convfun(void); +void nodupfun(void) __attribute__((noduplicate)); + +void f(void); +void g(void); + +// Test two if's are merged and non_convfun duplicated. +// The LLVM IR is equivalent to: +// if (a) { +// f(); +// non_convfun(); +// g(); +// } else { +// non_convfun(); +// } +// +// CHECK: define spir_func void @test_merge_if(i32 %[[a:.+]]) +// CHECK: %[[tobool:.+]] = icmp eq i32 %[[a]], 0 +// CHECK: br i1 %[[tobool]], label %[[if_end3_critedge:.+]], label %[[if_then:.+]] +// CHECK: [[if_then]]: +// CHECK: tail call spir_func void @f() +// CHECK: tail call spir_func void @non_convfun() +// CHECK: tail call spir_func void @g() +// CHECK: br label %[[if_end3:.+]] +// CHECK: [[if_end3_critedge]]: +// CHECK: tail call spir_func void @non_convfun() +// CHECK: br label %[[if_end3]] +// CHECK: [[if_end3]]: +// CHECK-LABEL: ret void + +void test_merge_if(int a) { + if (a) { + f(); + } + non_convfun(); + if (a) { + g(); + } +} + +// CHECK-DAG: declare spir_func void @f() +// CHECK-DAG: declare spir_func void @non_convfun() +// CHECK-DAG: declare spir_func void @g() + +// Test two if's are not merged. +// CHECK: define spir_func void @test_no_merge_if(i32 %[[a:.+]]) +// CHECK: %[[tobool:.+]] = icmp eq i32 %[[a]], 0 +// CHECK: br i1 %[[tobool]], label %[[if_end:.+]], label %[[if_then:.+]] +// CHECK: [[if_then]]: +// CHECK: tail call spir_func void @f() +// CHECK-NOT: call spir_func void @convfun() +// CHECK-NOT: call spir_func void @g() +// CHECK: br label %[[if_end]] +// CHECK: [[if_end]]: +// CHECK: %[[tobool_pr:.+]] = phi i1 [ true, %[[if_then]] ], [ false, %{{.+}} ] +// CHECK: tail call spir_func void @convfun() #[[attr5:.+]] +// CHECK: br i1 %[[tobool_pr]], label %[[if_then2:.+]], label %[[if_end3:.+]] +// CHECK: [[if_then2]]: +// CHECK: tail call spir_func void @g() +// CHECK: br label %[[if_end3:.+]] +// CHECK: [[if_end3]]: +// CHECK-LABEL: ret void + +void test_no_merge_if(int a) { + if (a) { + f(); + } + convfun(); + if(a) { + g(); + } +} + +// CHECK: declare spir_func void @convfun(){{[^#]*}} #[[attr2:[0-9]+]] + +// Test loop is unrolled for convergent function. +// CHECK-LABEL: define spir_func void @test_unroll() +// CHECK: tail call spir_func void @convfun() #[[attr5:[0-9]+]] +// CHECK: tail call spir_func void @convfun() #[[attr5]] +// CHECK: tail call spir_func void @convfun() #[[attr5]] +// CHECK: tail call spir_func void @convfun() #[[attr5]] +// CHECK: tail call spir_func void @convfun() #[[attr5]] +// CHECK: tail call spir_func void @convfun() #[[attr5]] +// CHECK: tail call spir_func void @convfun() #[[attr5]] +// CHECK: tail call spir_func void @convfun() #[[attr5]] +// CHECK: tail call spir_func void @convfun() #[[attr5]] +// CHECK: tail call spir_func void @convfun() #[[attr5]] +// CHECK-LABEL: ret void + +void test_unroll() { + for (int i = 0; i < 10; i++) + convfun(); +} + +// Test loop is not unrolled for noduplicate function. +// CHECK-LABEL: define spir_func void @test_not_unroll() +// CHECK: br label %[[for_body:.+]] +// CHECK: [[for_cond_cleanup:.+]]: +// CHECK: ret void +// CHECK: [[for_body]]: +// CHECK: tail call spir_func void @nodupfun() #[[attr6:[0-9]+]] +// CHECK-NOT: call spir_func void @nodupfun() +// CHECK: br i1 %{{.+}}, label %[[for_body]], label %[[for_cond_cleanup]] + +void test_not_unroll() { + for (int i = 0; i < 10; i++) + nodupfun(); +} + +// CHECK: declare spir_func void @nodupfun(){{[^#]*}} #[[attr3:[0-9]+]] + +// CHECK-DAG: attributes #[[attr2]] = { {{[^}]*}}convergent{{[^}]*}} } +// CHECK-DAG: attributes #[[attr3]] = { {{[^}]*}}noduplicate{{[^}]*}} } +// CHECK-DAG: attributes #[[attr5]] = { {{[^}]*}}convergent{{[^}]*}} } +// CHECK-DAG: attributes #[[attr6]] = { {{[^}]*}}noduplicate{{[^}]*}} } diff --git a/test/CodeGenOpenCL/denorms-are-zero.cl b/test/CodeGenOpenCL/denorms-are-zero.cl index 488004fa4fc31..f398218b4fd78 100644 --- a/test/CodeGenOpenCL/denorms-are-zero.cl +++ b/test/CodeGenOpenCL/denorms-are-zero.cl @@ -1,5 +1,19 @@ // RUN: %clang_cc1 -S -cl-denorms-are-zero -o - %s 2>&1 +// RUN: %clang_cc1 -emit-llvm -cl-denorms-are-zero -o - -triple amdgcn--amdhsa -target-cpu fiji %s | FileCheck %s +// RUN: %clang_cc1 -emit-llvm -o - -triple amdgcn--amdhsa -target-cpu fiji %s | FileCheck %s --check-prefix=CHECK-DENORM +// RUN: %clang_cc1 -emit-llvm -target-feature +fp32-denormals -target-feature -fp64-denormals -cl-denorms-are-zero -o - -triple amdgcn--amdhsa -target-cpu fiji %s | FileCheck --check-prefix=CHECK-FEATURE %s -// This test just checks that the -cl-denorms-are-zero argument is accepted +// For non-amdgcn targets, this test just checks that the -cl-denorms-are-zero argument is accepted // by clang. This option is currently a no-op, which is allowed by the // OpenCL specification. + +// For amdgcn target cpu fiji, fp32 should be flushed since fiji does not support fp32 denormals, unless +fp32-denormals is +// explicitly set. amdgcn target always do not flush fp64 denormals. + +// CHECK-DENORM-LABEL: define void @f() +// CHECK-DENORM: attributes #{{[0-9]*}} = {{{[^}]*}} "target-features"="{{[^"]*}}+fp64-denormals,{{[^"]*}}-fp32-denormals{{[^"]*}}" +// CHECK-LABEL: define void @f() +// CHECK: attributes #{{[0-9]*}} = {{{[^}]*}} "target-features"="{{[^"]*}}+fp64-denormals,{{[^"]*}}-fp32-denormals{{[^"]*}}" +// CHECK-FEATURE-LABEL: define void @f() +// CHECK-FEATURE: attributes #{{[0-9]*}} = {{{[^}]*}} "target-features"="{{[^"]*}}+fp32-denormals,{{[^"]*}}-fp64-denormals{{[^"]*}}" +void f() {} diff --git a/test/CodeGenOpenCL/extension-begin.cl b/test/CodeGenOpenCL/extension-begin.cl new file mode 100644 index 0000000000000..1d5f78965f34e --- /dev/null +++ b/test/CodeGenOpenCL/extension-begin.cl @@ -0,0 +1,25 @@ +// RUN: %clang_cc1 %s -triple spir-unknown-unknown -emit-llvm -o - | FileCheck %s + +__attribute__((overloadable)) void f(int x); + +#pragma OPENCL EXTENSION my_ext : begin + +__attribute__((overloadable)) void f(long x); + +#pragma OPENCL EXTENSION my_ext : end + +#pragma OPENCL EXTENSION my_ext : enable + +//CHECK: define spir_func void @test_f1(i64 %x) +//CHECK: call spir_func void @_Z1fl(i64 %{{.*}}) +void test_f1(long x) { + f(x); +} + +#pragma OPENCL EXTENSION my_ext : disable + +//CHECK: define spir_func void @test_f2(i64 %x) +//CHECK: call spir_func void @_Z1fi(i32 %{{.*}}) +void test_f2(long x) { + f(x); +} diff --git a/test/CodeGenOpenCL/fpmath.cl b/test/CodeGenOpenCL/fpmath.cl index 88df3bf166abb..8908861ace8e7 100644 --- a/test/CodeGenOpenCL/fpmath.cl +++ b/test/CodeGenOpenCL/fpmath.cl @@ -1,25 +1,51 @@ -// RUN: %clang_cc1 %s -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s +// RUN: %clang_cc1 %s -emit-llvm -o - -triple spir-unknown-unknown | FileCheck --check-prefix=CHECK --check-prefix=NODIVOPT %s +// RUN: %clang_cc1 %s -emit-llvm -o - -triple spir-unknown-unknown -cl-fp32-correctly-rounded-divide-sqrt | FileCheck --check-prefix=CHECK --check-prefix=DIVOPT %s +// RUN: %clang_cc1 %s -emit-llvm -o - -DNOFP64 -cl-std=CL1.2 -triple r600-unknown-unknown -target-cpu r600 -pedantic | FileCheck --check-prefix=CHECK-FLT %s +// RUN: %clang_cc1 %s -emit-llvm -o - -DFP64 -cl-std=CL1.2 -triple spir-unknown-unknown -pedantic | FileCheck --check-prefix=CHECK-DBL %s typedef __attribute__(( ext_vector_type(4) )) float float4; float spscalardiv(float a, float b) { // CHECK: @spscalardiv - // CHECK: fdiv{{.*}}, !fpmath ![[MD:[0-9]+]] + // CHECK: #[[ATTR:[0-9]+]] + // CHECK: fdiv{{.*}}, + // NODIVOPT: !fpmath ![[MD:[0-9]+]] + // DIVOPT-NOT: !fpmath ![[MD:[0-9]+]] return a / b; } float4 spvectordiv(float4 a, float4 b) { // CHECK: @spvectordiv - // CHECK: fdiv{{.*}}, !fpmath ![[MD]] + // CHECK: #[[ATTR]] + // CHECK: fdiv{{.*}}, + // NODIVOPT: !fpmath ![[MD]] + // DIVOPT-NOT: !fpmath ![[MD]] return a / b; } -#pragma OPENCL EXTENSION cl_khr_fp64 : enable +#if __OPENCL_C_VERSION__ >=120 +void printf(constant char* fmt, ...); + +void testdbllit(long *val) { + // CHECK-FLT: float 2.000000e+01 + // CHECK-DBL: double 2.000000e+01 + printf("%f", 20.0); +} +#endif + +#ifndef NOFP64 +#pragma OPENCL EXTENSION cl_khr_fp64 : enable double dpscalardiv(double a, double b) { // CHECK: @dpscalardiv + // CHECK: #[[ATTR]] // CHECK-NOT: !fpmath return a / b; } +#endif -// CHECK: ![[MD]] = !{float 2.500000e+00} +// CHECK: attributes #[[ATTR]] = { +// NODIVOPT: "correctly-rounded-divide-sqrt-fp-math"="false" +// DIVOPT: "correctly-rounded-divide-sqrt-fp-math"="true" +// CHECK: } +// NODIVOPT: ![[MD]] = !{float 2.500000e+00} diff --git a/test/CodeGenOpenCL/kernel-arg-info.cl b/test/CodeGenOpenCL/kernel-arg-info.cl index 5a5c8f9d6c5f6..32017f85df8cd 100644 --- a/test/CodeGenOpenCL/kernel-arg-info.cl +++ b/test/CodeGenOpenCL/kernel-arg-info.cl @@ -1,5 +1,5 @@ -// RUN: %clang_cc1 %s -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s -// RUN: %clang_cc1 %s -emit-llvm -o - -triple spir-unknown-unknown -cl-kernel-arg-info | FileCheck %s -check-prefix ARGINFO +// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s +// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -triple spir-unknown-unknown -cl-kernel-arg-info | FileCheck %s -check-prefix ARGINFO kernel void foo(__global int * restrict X, const int Y, volatile int anotherArg, __constant float * restrict Z) { @@ -14,7 +14,7 @@ kernel void foo(__global int * restrict X, const int Y, // CHECK-NOT: !kernel_arg_name // ARGINFO: !kernel_arg_name ![[MD15:[0-9]+]] -kernel void foo2(read_only image1d_t img1, image2d_t img2, write_only image2d_array_t img3) { +kernel void foo2(read_only image1d_t img1, image2d_t img2, write_only image2d_array_t img3, read_write image1d_t img4) { } // CHECK: define spir_kernel void @foo2{{[^!]+}} // CHECK: !kernel_arg_addr_space ![[MD21:[0-9]+]] @@ -65,11 +65,11 @@ kernel void foo5(myImage img1, write_only image1d_t img2) { // CHECK: ![[MD13]] = !{!"int*", !"int", !"int", !"float*"} // CHECK: ![[MD14]] = !{!"restrict", !"const", !"volatile", !"restrict const"} // ARGINFO: ![[MD15]] = !{!"X", !"Y", !"anotherArg", !"Z"} -// CHECK: ![[MD21]] = !{i32 1, i32 1, i32 1} -// CHECK: ![[MD22]] = !{!"read_only", !"read_only", !"write_only"} -// CHECK: ![[MD23]] = !{!"__read_only image1d_t", !"__read_only image2d_t", !"__write_only image2d_array_t"} -// CHECK: ![[MD24]] = !{!"", !"", !""} -// ARGINFO: ![[MD25]] = !{!"img1", !"img2", !"img3"} +// CHECK: ![[MD21]] = !{i32 1, i32 1, i32 1, i32 1} +// CHECK: ![[MD22]] = !{!"read_only", !"read_only", !"write_only", !"read_write"} +// CHECK: ![[MD23]] = !{!"image1d_t", !"image2d_t", !"image2d_array_t", !"image1d_t"} +// CHECK: ![[MD24]] = !{!"", !"", !"", !""} +// ARGINFO: ![[MD25]] = !{!"img1", !"img2", !"img3", !"img4"} // CHECK: ![[MD31]] = !{i32 1} // CHECK: ![[MD32]] = !{!"none"} // CHECK: ![[MD33]] = !{!"half*"} @@ -82,7 +82,7 @@ kernel void foo5(myImage img1, write_only image1d_t img2) { // CHECK: ![[MD45]] = !{!"", !""} // ARGINFO: ![[MD46]] = !{!"X", !"Y"} // CHECK: ![[MD51]] = !{!"read_only", !"write_only"} -// CHECK: ![[MD52]] = !{!"myImage", !"__write_only image1d_t"} -// CHECK: ![[MD53]] = !{!"__read_only image1d_t", !"__write_only image1d_t"} +// CHECK: ![[MD52]] = !{!"myImage", !"image1d_t"} +// CHECK: ![[MD53]] = !{!"image1d_t", !"image1d_t"} // ARGINFO: ![[MD54]] = !{!"img1", !"img2"} diff --git a/test/CodeGenOpenCL/local-initializer-undef.cl b/test/CodeGenOpenCL/local-initializer-undef.cl index 5d34f56b821d1..f209b1f6b449e 100644 --- a/test/CodeGenOpenCL/local-initializer-undef.cl +++ b/test/CodeGenOpenCL/local-initializer-undef.cl @@ -6,10 +6,10 @@ typedef struct Foo { float z; } Foo; -// CHECK-DAG: @test.lds_int = internal addrspace(2) global i32 undef -// CHECK-DAG: @test.lds_int_arr = internal addrspace(2) global [128 x i32] undef -// CHECK-DAG: @test.lds_struct = internal addrspace(2) global %struct.Foo undef -// CHECK-DAG: @test.lds_struct_arr = internal addrspace(2) global [64 x %struct.Foo] undef +// CHECK-DAG: @test.lds_int = internal addrspace(3) global i32 undef +// CHECK-DAG: @test.lds_int_arr = internal addrspace(3) global [128 x i32] undef +// CHECK-DAG: @test.lds_struct = internal addrspace(3) global %struct.Foo undef +// CHECK-DAG: @test.lds_struct_arr = internal addrspace(3) global [64 x %struct.Foo] undef __kernel void test() { __local int lds_int; diff --git a/test/CodeGenOpenCL/local.cl b/test/CodeGenOpenCL/local.cl index da371f8e46485..6f44b684e94bd 100644 --- a/test/CodeGenOpenCL/local.cl +++ b/test/CodeGenOpenCL/local.cl @@ -3,7 +3,7 @@ void func(local int*); __kernel void foo(void) { - // CHECK: @foo.i = internal addrspace(2) global i32 undef + // CHECK: @foo.i = internal addrspace(3) global i32 undef __local int i; func(&i); } diff --git a/test/CodeGenOpenCL/memcpy.cl b/test/CodeGenOpenCL/memcpy.cl index cb27803adf418..5911b5c5aa630 100644 --- a/test/CodeGenOpenCL/memcpy.cl +++ b/test/CodeGenOpenCL/memcpy.cl @@ -2,7 +2,7 @@ // CHECK-LABEL: @test // CHECK-NOT: addrspacecast -// CHECK: call void @llvm.memcpy.p1i8.p3i8 +// CHECK: call void @llvm.memcpy.p1i8.p2i8 kernel void test(global float *g, constant float *c) { __builtin_memcpy(g, c, 32); } diff --git a/test/CodeGenOpenCL/null_queue.cl b/test/CodeGenOpenCL/null_queue.cl new file mode 100644 index 0000000000000..cdcd7eef2ed46 --- /dev/null +++ b/test/CodeGenOpenCL/null_queue.cl @@ -0,0 +1,18 @@ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -emit-llvm %s -o - | FileCheck %s +extern queue_t get_default_queue(); + +bool compare() { + return 0 == get_default_queue() && + get_default_queue() == 0; + // CHECK: icmp eq %opencl.queue_t* null, %{{.*}} + // CHECK: icmp eq %opencl.queue_t* %{{.*}}, null +} + +void func(queue_t q); + +void init() { + queue_t q = 0; + func(0); + // CHECK: store %opencl.queue_t* null, %opencl.queue_t** %q + // CHECK: call void @func(%opencl.queue_t* null) +} diff --git a/test/CodeGenOpenCL/opencl_types.cl b/test/CodeGenOpenCL/opencl_types.cl index 5691b27fbe526..73c57b73098f2 100644 --- a/test/CodeGenOpenCL/opencl_types.cl +++ b/test/CodeGenOpenCL/opencl_types.cl @@ -1,40 +1,54 @@ -// RUN: %clang_cc1 %s -emit-llvm -o - -O0 | FileCheck %s +// RUN: %clang_cc1 %s -triple "spir-unknown-unknown" -emit-llvm -o - -O0 | FileCheck %s --check-prefix=CHECK-SPIR +// RUN: %clang_cc1 %s -triple "amdgcn--amdhsa" -emit-llvm -o - -O0 | FileCheck %s --check-prefix=CHECK-AMDGCN -constant sampler_t glb_smp = 7; -// CHECK: constant i32 7 +#define CLK_ADDRESS_CLAMP_TO_EDGE 2 +#define CLK_NORMALIZED_COORDS_TRUE 1 +#define CLK_FILTER_NEAREST 0x10 +#define CLK_FILTER_LINEAR 0x20 + +constant sampler_t glb_smp = CLK_ADDRESS_CLAMP_TO_EDGE|CLK_NORMALIZED_COORDS_TRUE|CLK_FILTER_NEAREST; +// CHECK-SPIR-NOT: constant i32 void fnc1(image1d_t img) {} -// CHECK: @fnc1(%opencl.image1d_ro_t* +// CHECK-SPIR: @fnc1(%opencl.image1d_ro_t addrspace(1)* +// CHECK-AMDGCN: @fnc1(%opencl.image1d_ro_t addrspace(2)* void fnc1arr(image1d_array_t img) {} -// CHECK: @fnc1arr(%opencl.image1d_array_ro_t* +// CHECK-SPIR: @fnc1arr(%opencl.image1d_array_ro_t addrspace(1)* +// CHECK-AMDGCN: @fnc1arr(%opencl.image1d_array_ro_t addrspace(2)* void fnc1buff(image1d_buffer_t img) {} -// CHECK: @fnc1buff(%opencl.image1d_buffer_ro_t* +// CHECK-SPIR: @fnc1buff(%opencl.image1d_buffer_ro_t addrspace(1)* +// CHECK-AMDGCN: @fnc1buff(%opencl.image1d_buffer_ro_t addrspace(2)* void fnc2(image2d_t img) {} -// CHECK: @fnc2(%opencl.image2d_ro_t* +// CHECK-SPIR: @fnc2(%opencl.image2d_ro_t addrspace(1)* +// CHECK-AMDGCN: @fnc2(%opencl.image2d_ro_t addrspace(2)* void fnc2arr(image2d_array_t img) {} -// CHECK: @fnc2arr(%opencl.image2d_array_ro_t* +// CHECK-SPIR: @fnc2arr(%opencl.image2d_array_ro_t addrspace(1)* +// CHECK-AMDGCN: @fnc2arr(%opencl.image2d_array_ro_t addrspace(2)* void fnc3(image3d_t img) {} -// CHECK: @fnc3(%opencl.image3d_ro_t* +// CHECK-SPIR: @fnc3(%opencl.image3d_ro_t addrspace(1)* +// CHECK-AMDGCN: @fnc3(%opencl.image3d_ro_t addrspace(2)* void fnc4smp(sampler_t s) {} -// CHECK-LABEL: define {{.*}}void @fnc4smp(i32 +// CHECK-SPIR-LABEL: define {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(2)* +// CHECK-AMDGCN-LABEL: define {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(2)* kernel void foo(image1d_t img) { - sampler_t smp = 5; - // CHECK: alloca i32 + sampler_t smp = CLK_ADDRESS_CLAMP_TO_EDGE|CLK_NORMALIZED_COORDS_TRUE|CLK_FILTER_LINEAR; + // CHECK-SPIR: alloca %opencl.sampler_t addrspace(2)* event_t evt; - // CHECK: alloca %opencl.event_t* - // CHECK: store i32 5, + // CHECK-SPIR: alloca %opencl.event_t* + // CHECK-SPIR: store %opencl.sampler_t addrspace(2)* fnc4smp(smp); - // CHECK: call {{.*}}void @fnc4smp(i32 + // CHECK-SPIR: call {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(2)* fnc4smp(glb_smp); - // CHECK: call {{.*}}void @fnc4smp(i32 + // CHECK-SPIR: call {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(2)* } void __attribute__((overloadable)) bad1(image1d_t b, image2d_t c, image2d_t d) {} -// CHECK-LABEL: @{{_Z4bad114ocl_image1d_ro14ocl_image2d_roS0_|"\\01\?bad1@@\$\$J0YAXPAUocl_image1d_ro@@PAUocl_image2d_ro@@1@Z"}} +// CHECK-SPIR-LABEL: @{{_Z4bad114ocl_image1d_ro14ocl_image2d_roS0_|"\\01\?bad1@@\$\$J0YAXPAUocl_image1d_ro@@PAUocl_image2d_ro@@1@Z"}} +// CHECK-AMDGCN-LABEL: @{{_Z4bad114ocl_image1d_ro14ocl_image2d_roS0_|"\\01\?bad1@@\$\$J0YAXPAUocl_image1d_ro@@PAUocl_image2d_ro@@1@Z"}}(%opencl.image1d_ro_t addrspace(2)*{{.*}}%opencl.image2d_ro_t addrspace(2)*{{.*}}%opencl.image2d_ro_t addrspace(2)*{{.*}}) diff --git a/test/CodeGenOpenCL/partial_initializer.cl b/test/CodeGenOpenCL/partial_initializer.cl new file mode 100644 index 0000000000000..454c82f892993 --- /dev/null +++ b/test/CodeGenOpenCL/partial_initializer.cl @@ -0,0 +1,66 @@ +// RUN: %clang_cc1 -triple spir-unknown-unknown -cl-std=CL2.0 -emit-llvm %s -O0 -o - | FileCheck %s + +typedef __attribute__(( ext_vector_type(2) )) int int2; +typedef __attribute__(( ext_vector_type(4) )) int int4; + +// CHECK: %struct.StrucTy = type { i32, i32, i32 } + +// CHECK: @GA = addrspace(1) global [6 x [6 x float]] {{[[][[]}}6 x float] [float 1.000000e+00, float 2.000000e+00, float 0.000000e+00, float 0.000000e+00, float 0.000000e+00, float 0.000000e+00], +// CHECK: [6 x float] zeroinitializer, [6 x float] zeroinitializer, [6 x float] zeroinitializer, [6 x float] zeroinitializer, [6 x float] zeroinitializer], align 4 +float GA[6][6] = {1.0f, 2.0f}; + +typedef struct { + int x; + int y; + int z; +} StrucTy; + +// CHECK: @GS = addrspace(1) global %struct.StrucTy { i32 1, i32 2, i32 0 }, align 4 +StrucTy GS = {1, 2}; + +// CHECK: @GV1 = addrspace(1) global <4 x i32> <i32 1, i32 2, i32 3, i32 4>, align 16 +int4 GV1 = (int4)((int2)(1,2),3,4); + +// CHECK: @GV2 = addrspace(1) global <4 x i32> <i32 1, i32 1, i32 1, i32 1>, align 16 +int4 GV2 = (int4)(1); + +// CHECK: @f.S = private unnamed_addr addrspace(2) constant %struct.StrucTy { i32 1, i32 2, i32 0 }, align 4 + +// CHECK-LABEL: define spir_func void @f() +void f(void) { + // CHECK: %[[A:.*]] = alloca [6 x [6 x float]], align 4 + // CHECK: %[[S:.*]] = alloca %struct.StrucTy, align 4 + // CHECK: %[[V1:.*]] = alloca <4 x i32>, align 16 + // CHECK: %[[compoundliteral:.*]] = alloca <4 x i32>, align 16 + // CHECK: %[[compoundliteral1:.*]] = alloca <2 x i32>, align 8 + // CHECK: %[[V2:.*]] = alloca <4 x i32>, align 16 + + // CHECK: %[[v0:.*]] = bitcast [6 x [6 x float]]* %A to i8* + // CHECK: call void @llvm.memset.p0i8.i32(i8* %[[v0]], i8 0, i32 144, i32 4, i1 false) + // CHECK: %[[v1:.*]] = bitcast i8* %[[v0]] to [6 x [6 x float]]* + // CHECK: %[[v2:.*]] = getelementptr [6 x [6 x float]], [6 x [6 x float]]* %[[v1]], i32 0, i32 0 + // CHECK: %[[v3:.*]] = getelementptr [6 x float], [6 x float]* %[[v2]], i32 0, i32 0 + // CHECK: store float 1.000000e+00, float* %[[v3]] + // CHECK: %[[v4:.*]] = getelementptr [6 x float], [6 x float]* %[[v2]], i32 0, i32 1 + // CHECK: store float 2.000000e+00, float* %[[v4]] + float A[6][6] = {1.0f, 2.0f}; + + // CHECK: %[[v5:.*]] = bitcast %struct.StrucTy* %S to i8* + // CHECK: call void @llvm.memcpy.p0i8.p2i8.i32(i8* %[[v5]], i8 addrspace(2)* bitcast (%struct.StrucTy addrspace(2)* @f.S to i8 addrspace(2)*), i32 12, i32 4, i1 false) + StrucTy S = {1, 2}; + + // CHECK: store <2 x i32> <i32 1, i32 2>, <2 x i32>* %[[compoundliteral1]], align 8 + // CHECK: %[[v6:.*]] = load <2 x i32>, <2 x i32>* %[[compoundliteral1]], align 8 + // CHECK: %[[vext:.*]] = shufflevector <2 x i32> %[[v6]], <2 x i32> undef, <4 x i32> <i32 0, i32 1, i32 undef, i32 undef> + // CHECK: %[[vecinit:.*]] = shufflevector <4 x i32> %[[vext]], <4 x i32> undef, <4 x i32> <i32 0, i32 1, i32 undef, i32 undef> + // CHECK: %[[vecinit2:.*]] = insertelement <4 x i32> %[[vecinit]], i32 3, i32 2 + // CHECK: %[[vecinit3:.*]] = insertelement <4 x i32> %[[vecinit2]], i32 4, i32 3 + // CHECK: store <4 x i32> %[[vecinit3]], <4 x i32>* %[[compoundliteral]], align 16 + // CHECK: %[[v7:.*]] = load <4 x i32>, <4 x i32>* %[[compoundliteral]], align 16 + // CHECK: store <4 x i32> %[[v7]], <4 x i32>* %[[V1]], align 16 + int4 V1 = (int4)((int2)(1,2),3,4); + + // CHECK: store <4 x i32> <i32 1, i32 1, i32 1, i32 1>, <4 x i32>* %[[V2]], align 16 + int4 V2 = (int4)(1); +} + diff --git a/test/CodeGenOpenCL/pipe_builtin.cl b/test/CodeGenOpenCL/pipe_builtin.cl index db6893ebf55d6..a9b4ab630cef8 100644 --- a/test/CodeGenOpenCL/pipe_builtin.cl +++ b/test/CodeGenOpenCL/pipe_builtin.cl @@ -4,58 +4,74 @@ // CHECK: %opencl.reserve_id_t = type opaque void test1(read_only pipe int p, global int *ptr) { - // CHECK: call i32 @__read_pipe_2(%opencl.pipe_t* %{{.*}}, i8* %{{.*}}) + // CHECK: call i32 @__read_pipe_2(%opencl.pipe_t* %{{.*}}, i8* %{{.*}}, i32 4, i32 4) read_pipe(p, ptr); - // CHECK: call %opencl.reserve_id_t* @__reserve_read_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}}) + // CHECK: call %opencl.reserve_id_t* @__reserve_read_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4) reserve_id_t rid = reserve_read_pipe(p, 2); - // CHECK: call i32 @__read_pipe_4(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 {{.*}}, i8* %{{.*}}) + // CHECK: call i32 @__read_pipe_4(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 {{.*}}, i8* %{{.*}}, i32 4, i32 4) read_pipe(p, rid, 2, ptr); - // CHECK: call void @__commit_read_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}) + // CHECK: call void @__commit_read_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4) commit_read_pipe(p, rid); } void test2(write_only pipe int p, global int *ptr) { - // CHECK: call i32 @__write_pipe_2(%opencl.pipe_t* %{{.*}}, i8* %{{.*}}) + // CHECK: call i32 @__write_pipe_2(%opencl.pipe_t* %{{.*}}, i8* %{{.*}}, i32 4, i32 4) write_pipe(p, ptr); - // CHECK: call %opencl.reserve_id_t* @__reserve_write_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}}) + // CHECK: call %opencl.reserve_id_t* @__reserve_write_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4) reserve_id_t rid = reserve_write_pipe(p, 2); - // CHECK: call i32 @__write_pipe_4(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 {{.*}}, i8* %{{.*}}) + // CHECK: call i32 @__write_pipe_4(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 {{.*}}, i8* %{{.*}}, i32 4, i32 4) write_pipe(p, rid, 2, ptr); - // CHECK: call void @__commit_write_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}) + // CHECK: call void @__commit_write_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4) commit_write_pipe(p, rid); } void test3(read_only pipe int p, global int *ptr) { - // CHECK: call %opencl.reserve_id_t* @__work_group_reserve_read_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}}) + // CHECK: call %opencl.reserve_id_t* @__work_group_reserve_read_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4) reserve_id_t rid = work_group_reserve_read_pipe(p, 2); - // CHECK: call void @__work_group_commit_read_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}) + // CHECK: call void @__work_group_commit_read_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4) work_group_commit_read_pipe(p, rid); } void test4(write_only pipe int p, global int *ptr) { - // CHECK: call %opencl.reserve_id_t* @__work_group_reserve_write_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}}) + // CHECK: call %opencl.reserve_id_t* @__work_group_reserve_write_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4) reserve_id_t rid = work_group_reserve_write_pipe(p, 2); - // CHECK: call void @__work_group_commit_write_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}) + // CHECK: call void @__work_group_commit_write_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4) work_group_commit_write_pipe(p, rid); } void test5(read_only pipe int p, global int *ptr) { - // CHECK: call %opencl.reserve_id_t* @__sub_group_reserve_read_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}}) + // CHECK: call %opencl.reserve_id_t* @__sub_group_reserve_read_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4) reserve_id_t rid = sub_group_reserve_read_pipe(p, 2); - // CHECK: call void @__sub_group_commit_read_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}) + // CHECK: call void @__sub_group_commit_read_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4) sub_group_commit_read_pipe(p, rid); } void test6(write_only pipe int p, global int *ptr) { - // CHECK: call %opencl.reserve_id_t* @__sub_group_reserve_write_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}}) + // CHECK: call %opencl.reserve_id_t* @__sub_group_reserve_write_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4) reserve_id_t rid = sub_group_reserve_write_pipe(p, 2); - // CHECK: call void @__sub_group_commit_write_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}) + // CHECK: call void @__sub_group_commit_write_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4) sub_group_commit_write_pipe(p, rid); } void test7(write_only pipe int p, global int *ptr) { - // CHECK: call i32 @__get_pipe_num_packets(%opencl.pipe_t* %{{.*}}) + // CHECK: call i32 @__get_pipe_num_packets(%opencl.pipe_t* %{{.*}}, i32 4, i32 4) *ptr = get_pipe_num_packets(p); - // CHECK: call i32 @__get_pipe_max_packets(%opencl.pipe_t* %{{.*}}) + // CHECK: call i32 @__get_pipe_max_packets(%opencl.pipe_t* %{{.*}}, i32 4, i32 4) *ptr = get_pipe_max_packets(p); } + +void test8(read_only pipe int r, write_only pipe int w, global int *ptr) { + // verify that return type is correctly casted to i1 value + // CHECK: %[[R:[0-9]+]] = call i32 @__read_pipe_2 + // CHECK: icmp ne i32 %[[R]], 0 + if (read_pipe(r, ptr)) *ptr = -1; + // CHECK: %[[W:[0-9]+]] = call i32 @__write_pipe_2 + // CHECK: icmp ne i32 %[[W]], 0 + if (write_pipe(w, ptr)) *ptr = -1; + // CHECK: %[[N:[0-9]+]] = call i32 @__get_pipe_num_packets + // CHECK: icmp ne i32 %[[N]], 0 + if (get_pipe_num_packets(r)) *ptr = -1; + // CHECK: %[[M:[0-9]+]] = call i32 @__get_pipe_max_packets + // CHECK: icmp ne i32 %[[M]], 0 + if (get_pipe_max_packets(w)) *ptr = -1; +} diff --git a/test/CodeGenOpenCL/pipe_types.cl b/test/CodeGenOpenCL/pipe_types.cl index b9c411b5fcb6a..f5b42e2a490fc 100644 --- a/test/CodeGenOpenCL/pipe_types.cl +++ b/test/CodeGenOpenCL/pipe_types.cl @@ -41,7 +41,7 @@ void test_reserved_read_pipe(global struct Person *SDst, read_only pipe struct Person SPipe) { // CHECK: define void @test_reserved_read_pipe read_pipe (SPipe, SDst); - // CHECK: call i32 @__read_pipe_2(%opencl.pipe_t* %{{.*}}, i8* %{{.*}}) + // CHECK: call i32 @__read_pipe_2(%opencl.pipe_t* %{{.*}}, i8* %{{.*}}, i32 16, i32 8) read_pipe (SPipe, SDst); - // CHECK: call i32 @__read_pipe_2(%opencl.pipe_t* %{{.*}}, i8* %{{.*}}) + // CHECK: call i32 @__read_pipe_2(%opencl.pipe_t* %{{.*}}, i8* %{{.*}}, i32 16, i32 8) } diff --git a/test/CodeGenOpenCL/private-array-initialization.cl b/test/CodeGenOpenCL/private-array-initialization.cl new file mode 100644 index 0000000000000..a7d4a98919779 --- /dev/null +++ b/test/CodeGenOpenCL/private-array-initialization.cl @@ -0,0 +1,9 @@ +// RUN: %clang_cc1 %s -triple spir-unknown-unknown -O0 -emit-llvm -o - | FileCheck %s + +// CHECK: @test.arr = private unnamed_addr addrspace(2) constant [3 x i32] [i32 1, i32 2, i32 3], align 4 + +void test() { + __private int arr[] = {1, 2, 3}; +// CHECK: %[[arr_i8_ptr:[0-9]+]] = bitcast [3 x i32]* %arr to i8* +// CHECK: call void @llvm.memcpy.p0i8.p2i8.i32(i8* %[[arr_i8_ptr]], i8 addrspace(2)* bitcast ([3 x i32] addrspace(2)* @test.arr to i8 addrspace(2)*), i32 12, i32 4, i1 false) +} diff --git a/test/CodeGenOpenCL/sampler.cl b/test/CodeGenOpenCL/sampler.cl new file mode 100644 index 0000000000000..6fc8c2c684469 --- /dev/null +++ b/test/CodeGenOpenCL/sampler.cl @@ -0,0 +1,57 @@ +// RUN: %clang_cc1 %s -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck %s +// +// This test covers 5 cases of sampler initialzation: +// 1. function argument passing +// 1a. argument is a file-scope variable +// 1b. argument is a function-scope variable +// 1c. argument is one of caller function's parameters +// 2. variable initialization +// 2a. initializing a file-scope variable +// 2b. initializing a function-scope variable + +#define CLK_ADDRESS_CLAMP_TO_EDGE 2 +#define CLK_NORMALIZED_COORDS_TRUE 1 +#define CLK_FILTER_NEAREST 0x10 +#define CLK_FILTER_LINEAR 0x20 + +// CHECK: %opencl.sampler_t = type opaque + +// Case 2a +constant sampler_t glb_smp = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR; +// CHECK-NOT: glb_smp + +void fnc4smp(sampler_t s) {} +// CHECK: define spir_func void @fnc4smp(%opencl.sampler_t addrspace(2)* % + +kernel void foo(sampler_t smp_par) { + // CHECK-LABEL: define spir_kernel void @foo(%opencl.sampler_t addrspace(2)* %smp_par) + // CHECK: [[smp_par_ptr:%[A-Za-z0-9_\.]+]] = alloca %opencl.sampler_t addrspace(2)* + + // Case 2b + sampler_t smp = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_NEAREST; + // CHECK: [[smp_ptr:%[A-Za-z0-9_\.]+]] = alloca %opencl.sampler_t addrspace(2)* + // CHECK: [[SAMP:%[0-9]+]] = call %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 19) + // CHECK: store %opencl.sampler_t addrspace(2)* [[SAMP]], %opencl.sampler_t addrspace(2)** [[smp_ptr]] + + // Case 1b + fnc4smp(smp); + // CHECK-NOT: call %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 19) + // CHECK: [[SAMP:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[smp_ptr]] + // CHECK: call spir_func void @fnc4smp(%opencl.sampler_t addrspace(2)* [[SAMP]]) + + // Case 1b + fnc4smp(smp); + // CHECK-NOT: call %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 19) + // CHECK: [[SAMP:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[smp_ptr]] + // CHECK: call spir_func void @fnc4smp(%opencl.sampler_t addrspace(2)* [[SAMP]]) + + // Case 1a + fnc4smp(glb_smp); + // CHECK: [[SAMP:%[0-9]+]] = call %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 35) + // CHECK: call spir_func void @fnc4smp(%opencl.sampler_t addrspace(2)* [[SAMP]]) + + // Case 1c + fnc4smp(smp_par); + // CHECK: [[SAMP:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[smp_par_ptr]] + // CHECK: call spir_func void @fnc4smp(%opencl.sampler_t addrspace(2)* [[SAMP]]) +} diff --git a/test/CodeGenOpenCL/size_t.cl b/test/CodeGenOpenCL/size_t.cl new file mode 100644 index 0000000000000..ed941e01852d2 --- /dev/null +++ b/test/CodeGenOpenCL/size_t.cl @@ -0,0 +1,110 @@ +// RUN: %clang_cc1 %s -cl-std=CL2.0 -finclude-default-header -emit-llvm -O0 -triple spir-unknown-unknown -o - | FileCheck --check-prefix=SZ32 %s +// RUN: %clang_cc1 %s -cl-std=CL2.0 -finclude-default-header -emit-llvm -O0 -triple spir64-unknown-unknown -o - | FileCheck --check-prefix=SZ64 --check-prefix=SZ64ONLY %s +// RUN: %clang_cc1 %s -cl-std=CL2.0 -finclude-default-header -emit-llvm -O0 -triple amdgcn-- -o - | FileCheck --check-prefix=SZ64 --check-prefix=AMDONLY %s + +//SZ32: define{{.*}} i32 @test_ptrtoint_private(i8* %x) +//SZ32: ptrtoint i8* %{{.*}} to i32 +//SZ64: define{{.*}} i64 @test_ptrtoint_private(i8* %x) +//SZ64: ptrtoint i8* %{{.*}} to i64 +size_t test_ptrtoint_private(private char* x) { + return (size_t)x; +} + +//SZ32: define{{.*}} i32 @test_ptrtoint_global(i8 addrspace(1)* %x) +//SZ32: ptrtoint i8 addrspace(1)* %{{.*}} to i32 +//SZ64: define{{.*}} i64 @test_ptrtoint_global(i8 addrspace(1)* %x) +//SZ64: ptrtoint i8 addrspace(1)* %{{.*}} to i64 +intptr_t test_ptrtoint_global(global char* x) { + return (intptr_t)x; +} + +//SZ32: define{{.*}} i32 @test_ptrtoint_constant(i8 addrspace(2)* %x) +//SZ32: ptrtoint i8 addrspace(2)* %{{.*}} to i32 +//SZ64: define{{.*}} i64 @test_ptrtoint_constant(i8 addrspace(2)* %x) +//SZ64: ptrtoint i8 addrspace(2)* %{{.*}} to i64 +uintptr_t test_ptrtoint_constant(constant char* x) { + return (uintptr_t)x; +} + +//SZ32: define{{.*}} i32 @test_ptrtoint_local(i8 addrspace(3)* %x) +//SZ32: ptrtoint i8 addrspace(3)* %{{.*}} to i32 +//SZ64: define{{.*}} i64 @test_ptrtoint_local(i8 addrspace(3)* %x) +//SZ64: ptrtoint i8 addrspace(3)* %{{.*}} to i64 +size_t test_ptrtoint_local(local char* x) { + return (size_t)x; +} + +//SZ32: define{{.*}} i32 @test_ptrtoint_generic(i8 addrspace(4)* %x) +//SZ32: ptrtoint i8 addrspace(4)* %{{.*}} to i32 +//SZ64: define{{.*}} i64 @test_ptrtoint_generic(i8 addrspace(4)* %x) +//SZ64: ptrtoint i8 addrspace(4)* %{{.*}} to i64 +size_t test_ptrtoint_generic(generic char* x) { + return (size_t)x; +} + +//SZ32: define{{.*}} i8* @test_inttoptr_private(i32 %x) +//SZ32: inttoptr i32 %{{.*}} to i8* +//SZ64: define{{.*}} i8* @test_inttoptr_private(i64 %x) +//AMDONLY: trunc i64 %{{.*}} to i32 +//AMDONLY: inttoptr i32 %{{.*}} to i8* +//SZ64ONLY: inttoptr i64 %{{.*}} to i8* +private char* test_inttoptr_private(size_t x) { + return (private char*)x; +} + +//SZ32: define{{.*}} i8 addrspace(1)* @test_inttoptr_global(i32 %x) +//SZ32: inttoptr i32 %{{.*}} to i8 addrspace(1)* +//SZ64: define{{.*}} i8 addrspace(1)* @test_inttoptr_global(i64 %x) +//SZ64: inttoptr i64 %{{.*}} to i8 addrspace(1)* +global char* test_inttoptr_global(size_t x) { + return (global char*)x; +} + +//SZ32: define{{.*}} i8 addrspace(3)* @test_add_local(i8 addrspace(3)* %x, i32 %y) +//SZ32: getelementptr inbounds i8, i8 addrspace(3)* %{{.*}}, i32 +//SZ64: define{{.*}} i8 addrspace(3)* @test_add_local(i8 addrspace(3)* %x, i64 %y) +//AMDONLY: trunc i64 %{{.*}} to i32 +//AMDONLY: getelementptr inbounds i8, i8 addrspace(3)* %{{.*}}, i32 +//SZ64ONLY: getelementptr inbounds i8, i8 addrspace(3)* %{{.*}}, i64 +local char* test_add_local(local char* x, ptrdiff_t y) { + return x + y; +} + +//SZ32: define{{.*}} i8 addrspace(1)* @test_add_global(i8 addrspace(1)* %x, i32 %y) +//SZ32: getelementptr inbounds i8, i8 addrspace(1)* %{{.*}}, i32 +//SZ64: define{{.*}} i8 addrspace(1)* @test_add_global(i8 addrspace(1)* %x, i64 %y) +//SZ64: getelementptr inbounds i8, i8 addrspace(1)* %{{.*}}, i64 +global char* test_add_global(global char* x, ptrdiff_t y) { + return x + y; +} + +//SZ32: define{{.*}} i32 @test_sub_local(i8 addrspace(3)* %x, i8 addrspace(3)* %y) +//SZ32: ptrtoint i8 addrspace(3)* %{{.*}} to i32 +//SZ32: ptrtoint i8 addrspace(3)* %{{.*}} to i32 +//SZ64: define{{.*}} i64 @test_sub_local(i8 addrspace(3)* %x, i8 addrspace(3)* %y) +//SZ64: ptrtoint i8 addrspace(3)* %{{.*}} to i64 +//SZ64: ptrtoint i8 addrspace(3)* %{{.*}} to i64 +ptrdiff_t test_sub_local(local char* x, local char *y) { + return x - y; +} + +//SZ32: define{{.*}} i32 @test_sub_private(i8* %x, i8* %y) +//SZ32: ptrtoint i8* %{{.*}} to i32 +//SZ32: ptrtoint i8* %{{.*}} to i32 +//SZ64: define{{.*}} i64 @test_sub_private(i8* %x, i8* %y) +//SZ64: ptrtoint i8* %{{.*}} to i64 +//SZ64: ptrtoint i8* %{{.*}} to i64 +ptrdiff_t test_sub_private(private char* x, private char *y) { + return x - y; +} + +//SZ32: define{{.*}} i32 @test_sub_mix(i8* %x, i8 addrspace(4)* %y) +//SZ32: ptrtoint i8* %{{.*}} to i32 +//SZ32: ptrtoint i8 addrspace(4)* %{{.*}} to i32 +//SZ64: define{{.*}} i64 @test_sub_mix(i8* %x, i8 addrspace(4)* %y) +//SZ64: ptrtoint i8* %{{.*}} to i64 +//SZ64: ptrtoint i8 addrspace(4)* %{{.*}} to i64 +ptrdiff_t test_sub_mix(private char* x, generic char *y) { + return x - y; +} + diff --git a/test/CodeGenOpenCL/spir_version.cl b/test/CodeGenOpenCL/spir_version.cl index 215b4d7a40238..8a191282b3c78 100644 --- a/test/CodeGenOpenCL/spir_version.cl +++ b/test/CodeGenOpenCL/spir_version.cl @@ -1,18 +1,34 @@ -// RUN: %clang_cc1 %s -triple "spir-unknown-unknown" -emit-llvm -o - | FileCheck %s --check-prefix=CL10 -// RUN: %clang_cc1 %s -triple "spir-unknown-unknown" -emit-llvm -o - -cl-std=CL1.2 | FileCheck %s --check-prefix=CL12 -// RUN: %clang_cc1 %s -triple "spir-unknown-unknown" -emit-llvm -o - -cl-std=CL2.0 | FileCheck %s --check-prefix=CL20 -// RUN: %clang_cc1 %s -triple "spir64-unknown-unknown" -emit-llvm -o - | FileCheck %s --check-prefix=CL10 -// RUN: %clang_cc1 %s -triple "spir64-unknown-unknown" -emit-llvm -o - -cl-std=CL1.2 | FileCheck %s --check-prefix=CL12 -// RUN: %clang_cc1 %s -triple "spir64-unknown-unknown" -emit-llvm -o - -cl-std=CL2.0 | FileCheck %s --check-prefix=CL20 +// RUN: %clang_cc1 %s -triple "spir-unknown-unknown" -emit-llvm -o - | FileCheck %s --check-prefix=CHECK-SPIR-CL10 +// RUN: %clang_cc1 %s -triple "spir-unknown-unknown" -emit-llvm -o - -cl-std=CL1.2 | FileCheck %s --check-prefix=CHECK-SPIR-CL12 +// RUN: %clang_cc1 %s -triple "spir-unknown-unknown" -emit-llvm -o - -cl-std=CL2.0 | FileCheck %s --check-prefix=CHECK-SPIR-CL20 +// RUN: %clang_cc1 %s -triple "spir64-unknown-unknown" -emit-llvm -o - | FileCheck %s --check-prefix=CHECK-SPIR-CL10 +// RUN: %clang_cc1 %s -triple "spir64-unknown-unknown" -emit-llvm -o - -cl-std=CL1.2 | FileCheck %s --check-prefix=CHECK-SPIR-CL12 +// RUN: %clang_cc1 %s -triple "spir64-unknown-unknown" -emit-llvm -o - -cl-std=CL2.0 | FileCheck %s --check-prefix=CHECK-SPIR-CL20 + +// RUN: %clang_cc1 %s -triple "amdgcn--amdhsa" -emit-llvm -o - | FileCheck %s --check-prefix=CHECK-AMDGCN-CL10 +// RUN: %clang_cc1 %s -triple "amdgcn--amdhsa" -emit-llvm -o - -cl-std=CL1.2 | FileCheck %s --check-prefix=CHECK-AMDGCN-CL12 +// RUN: %clang_cc1 %s -triple "amdgcn--amdhsa" -emit-llvm -o - -cl-std=CL2.0 | FileCheck %s --check-prefix=CHECK-AMDGCN-CL20 + kernel void foo() {} -// CL10: !opencl.spir.version = !{[[SPIR:![0-9]+]]} -// CL10: !opencl.ocl.version = !{[[OCL:![0-9]+]]} -// CL10: [[SPIR]] = !{i32 2, i32 0} -// CL10: [[OCL]] = !{i32 1, i32 0} -// CL12: !opencl.spir.version = !{[[SPIR:![0-9]+]]} -// CL12: !opencl.ocl.version = !{[[OCL:![0-9]+]]} -// CL12: [[SPIR]] = !{i32 2, i32 0} -// CL12: [[OCL]] = !{i32 1, i32 2} -// CL20: !opencl.spir.version = !{[[SPIR:![0-9]+]]} -// CL20: !opencl.ocl.version = !{[[SPIR:![0-9]+]]} -// CL20: [[SPIR]] = !{i32 2, i32 0} + +// CHECK-SPIR-CL10: !opencl.spir.version = !{[[SPIR:![0-9]+]]} +// CHECK-SPIR-CL10: !opencl.ocl.version = !{[[OCL:![0-9]+]]} +// CHECK-SPIR-CL10: [[SPIR]] = !{i32 1, i32 2} +// CHECK-SPIR-CL10: [[OCL]] = !{i32 1, i32 0} +// CHECK-SPIR-CL12: !opencl.spir.version = !{[[VER:![0-9]+]]} +// CHECK-SPIR-CL12: !opencl.ocl.version = !{[[VER]]} +// CHECK-SPIR-CL12: [[VER]] = !{i32 1, i32 2} + +// CHECK-SPIR-CL20: !opencl.spir.version = !{[[VER:![0-9]+]]} +// CHECK-SPIR-CL20: !opencl.ocl.version = !{[[VER]]} +// CHECK-SPIR-CL20: [[VER]] = !{i32 2, i32 0} + +// CHECK-AMDGCN-CL10-NOT: !opencl.spir.version +// CHECK-AMDGCN-CL10: !opencl.ocl.version = !{[[OCL:![0-9]+]]} +// CHECK-AMDGCN-CL10: [[OCL]] = !{i32 1, i32 0} +// CHECK-AMDGCN-CL12-NOT: !opencl.spir.version +// CHECK-AMDGCN-CL12: !opencl.ocl.version = !{[[OCL:![0-9]+]]} +// CHECK-AMDGCN-CL12: [[OCL]] = !{i32 1, i32 2} +// CHECK-AMDGCN-CL20-NOT: !opencl.spir.version +// CHECK-AMDGCN-CL20: !opencl.ocl.version = !{[[OCL:![0-9]+]]} +// CHECK-AMDGCN-CL20: [[OCL]] = !{i32 2, i32 0} diff --git a/test/CodeGenOpenCL/str_literals.cl b/test/CodeGenOpenCL/str_literals.cl index 1c0acd155dac3..8fe5cefc8d757 100644 --- a/test/CodeGenOpenCL/str_literals.cl +++ b/test/CodeGenOpenCL/str_literals.cl @@ -3,7 +3,7 @@ __constant char * __constant x = "hello world"; __constant char * __constant y = "hello world"; -// CHECK: unnamed_addr addrspace(3) constant -// CHECK-NOT: addrspace(3) unnamed_addr constant -// CHECK: @x = addrspace(3) constant i8 addrspace(3)* -// CHECK: @y = addrspace(3) constant i8 addrspace(3)* +// CHECK: unnamed_addr addrspace(2) constant +// CHECK-NOT: addrspace(2) unnamed_addr constant +// CHECK: @x = addrspace(2) constant i8 addrspace(2)* +// CHECK: @y = addrspace(2) constant i8 addrspace(2)* diff --git a/test/CodeGenOpenCL/unroll-hint.cl b/test/CodeGenOpenCL/unroll-hint.cl index a86762e02b548..6a9ba87a5eb01 100644 --- a/test/CodeGenOpenCL/unroll-hint.cl +++ b/test/CodeGenOpenCL/unroll-hint.cl @@ -1,4 +1,5 @@ // RUN: %clang_cc1 -emit-llvm -O0 -cl-std=CL2.0 -o - %s | FileCheck %s +// RUN: %clang_cc1 -emit-llvm -O0 -cl-std=CL1.2 -o - %s | FileCheck %s /*** for ***/ void for_count() |