diff options
Diffstat (limited to 'test/CodeGenOpenCL')
22 files changed, 1265 insertions, 111 deletions
diff --git a/test/CodeGenOpenCL/address-space-constant-initializers.cl b/test/CodeGenOpenCL/address-space-constant-initializers.cl index 45c64c2f716b..a03d939fae54 100644 --- a/test/CodeGenOpenCL/address-space-constant-initializers.cl +++ b/test/CodeGenOpenCL/address-space-constant-initializers.cl @@ -1,4 +1,6 @@ // RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple amdgcn-amd-amdhsa-opencl -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple amdgcn-amd-amdhsa-amdgizcl -emit-llvm -o - | FileCheck %s typedef struct { int i; diff --git a/test/CodeGenOpenCL/address-spaces.cl b/test/CodeGenOpenCL/address-spaces.cl index 3b34986d8936..7c665286547a 100644 --- a/test/CodeGenOpenCL/address-spaces.cl +++ b/test/CodeGenOpenCL/address-spaces.cl @@ -1,7 +1,12 @@ -// RUN: %clang_cc1 %s -O0 -ffake-address-space-map -emit-llvm -o - | FileCheck %s -// RUN: %clang_cc1 %s -O0 -DCL20 -cl-std=CL2.0 -ffake-address-space-map -emit-llvm -o - | FileCheck %s --check-prefix=CL20 +// RUN: %clang_cc1 %s -O0 -ffake-address-space-map -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,SPIR +// RUN: %clang_cc1 %s -O0 -DCL20 -cl-std=CL2.0 -ffake-address-space-map -emit-llvm -o - | FileCheck %s --check-prefixes=CL20,CL20SPIR +// RUN: %clang_cc1 %s -O0 -triple amdgcn-amd-amdhsa-opencl -emit-llvm -o - | FileCheck --check-prefixes=CHECK,SPIR %s +// RUN: %clang_cc1 %s -O0 -triple amdgcn-amd-amdhsa-opencl -DCL20 -cl-std=CL2.0 -emit-llvm -o - | FileCheck %s --check-prefixes=CL20,CL20SPIR +// RUN: %clang_cc1 %s -O0 -triple amdgcn-amd-amdhsa-amdgizcl -emit-llvm -o - | FileCheck %s -check-prefixes=CHECK,GIZ +// RUN: %clang_cc1 %s -O0 -triple amdgcn-amd-amdhsa-amdgizcl -DCL20 -cl-std=CL2.0 -emit-llvm -o - | FileCheck %s --check-prefixes=CL20,CL20GIZ -// CHECK: i32* %arg +// SPIR: i32* %arg +// GIZ: i32 addrspace(5)* %arg void f__p(__private int *arg) {} // CHECK: i32 addrspace(1)* %arg @@ -13,7 +18,8 @@ void f__l(__local int *arg) {} // CHECK: i32 addrspace(2)* %arg void f__c(__constant int *arg) {} -// CHECK: i32* %arg +// SPIR: i32* %arg +// GIZ: i32 addrspace(5)* %arg void fp(private int *arg) {} // CHECK: i32 addrspace(1)* %arg @@ -29,16 +35,21 @@ void fc(constant int *arg) {} int i; // CL20-DAG: @i = common addrspace(1) global i32 0 int *ptr; -// CL20-DAG: @ptr = common addrspace(1) global i32 addrspace(4)* null +// CL20SPIR-DAG: @ptr = common addrspace(1) global i32 addrspace(4)* null +// CL20GIZ-DAG: @ptr = common addrspace(1) global i32* null #endif -// CHECK: i32* %arg -// CL20-DAG: i32 addrspace(4)* %arg +// SPIR: i32* %arg +// GIZ: i32 addrspace(5)* %arg +// CL20SPIR-DAG: i32 addrspace(4)* %arg +// CL20GIZ-DAG: i32* %arg void f(int *arg) { int i; -// CHECK: %i = alloca i32, -// CL20-DAG: %i = alloca i32, +// SPIR: %i = alloca i32, +// GIZ: %i = alloca i32{{.*}}addrspace(5) +// CL20SPIR-DAG: %i = alloca i32, +// CL20GIZ-DAG: %i = alloca i32{{.*}}addrspace(5) #ifdef CL20 static int ii; diff --git a/test/CodeGenOpenCL/amdgcn-large-globals.cl b/test/CodeGenOpenCL/amdgcn-large-globals.cl new file mode 100644 index 000000000000..ea9ea618a488 --- /dev/null +++ b/test/CodeGenOpenCL/amdgcn-large-globals.cl @@ -0,0 +1,12 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s + +// CHECK: @One = common local_unnamed_addr addrspace(1) global [6442450944 x i8] zeroinitializer, align 1 +unsigned char One[6442450944]; +// CHECK: @Two = common local_unnamed_addr addrspace(1) global [6442450944 x i32] zeroinitializer, align 4 +global unsigned int Two[6442450944]; + +kernel void large_globals(unsigned int id) { + One[id] = id; + Two[id + 1] = id + 1; +} diff --git a/test/CodeGenOpenCL/amdgpu-alignment.cl b/test/CodeGenOpenCL/amdgpu-alignment.cl new file mode 100644 index 000000000000..35a8342eec01 --- /dev/null +++ b/test/CodeGenOpenCL/amdgpu-alignment.cl @@ -0,0 +1,522 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -S -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s + +#pragma OPENCL EXTENSION cl_khr_fp64 : enable +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +typedef char __attribute__((ext_vector_type(2))) char2; +typedef char __attribute__((ext_vector_type(3))) char3; +typedef char __attribute__((ext_vector_type(4))) char4; +typedef char __attribute__((ext_vector_type(8))) char8; +typedef char __attribute__((ext_vector_type(16))) char16; + +typedef short __attribute__((ext_vector_type(2))) short2; +typedef short __attribute__((ext_vector_type(3))) short3; +typedef short __attribute__((ext_vector_type(4))) short4; +typedef short __attribute__((ext_vector_type(8))) short8; +typedef short __attribute__((ext_vector_type(16))) short16; + +typedef int __attribute__((ext_vector_type(2))) int2; +typedef int __attribute__((ext_vector_type(3))) int3; +typedef int __attribute__((ext_vector_type(4))) int4; +typedef int __attribute__((ext_vector_type(8))) int8; +typedef int __attribute__((ext_vector_type(16))) int16; + +typedef long __attribute__((ext_vector_type(2))) long2; +typedef long __attribute__((ext_vector_type(3))) long3; +typedef long __attribute__((ext_vector_type(4))) long4; +typedef long __attribute__((ext_vector_type(8))) long8; +typedef long __attribute__((ext_vector_type(16))) long16; + +typedef half __attribute__((ext_vector_type(2))) half2; +typedef half __attribute__((ext_vector_type(3))) half3; +typedef half __attribute__((ext_vector_type(4))) half4; +typedef half __attribute__((ext_vector_type(8))) half8; +typedef half __attribute__((ext_vector_type(16))) half16; + +typedef float __attribute__((ext_vector_type(2))) float2; +typedef float __attribute__((ext_vector_type(3))) float3; +typedef float __attribute__((ext_vector_type(4))) float4; +typedef float __attribute__((ext_vector_type(8))) float8; +typedef float __attribute__((ext_vector_type(16))) float16; + +typedef double __attribute__((ext_vector_type(2))) double2; +typedef double __attribute__((ext_vector_type(3))) double3; +typedef double __attribute__((ext_vector_type(4))) double4; +typedef double __attribute__((ext_vector_type(8))) double8; +typedef double __attribute__((ext_vector_type(16))) double16; + +// CHECK: @local_memory_alignment_global.lds_i8 = internal addrspace(3) global [4 x i8] undef, align 1 +// CHECK: @local_memory_alignment_global.lds_v2i8 = internal addrspace(3) global [4 x <2 x i8>] undef, align 2 +// CHECK: @local_memory_alignment_global.lds_v3i8 = internal addrspace(3) global [4 x <3 x i8>] undef, align 4 +// CHECK: @local_memory_alignment_global.lds_v4i8 = internal addrspace(3) global [4 x <4 x i8>] undef, align 4 +// CHECK: @local_memory_alignment_global.lds_v8i8 = internal addrspace(3) global [4 x <8 x i8>] undef, align 8 +// CHECK: @local_memory_alignment_global.lds_v16i8 = internal addrspace(3) global [4 x <16 x i8>] undef, align 16 +// CHECK: @local_memory_alignment_global.lds_i16 = internal addrspace(3) global [4 x i16] undef, align 2 +// CHECK: @local_memory_alignment_global.lds_v2i16 = internal addrspace(3) global [4 x <2 x i16>] undef, align 4 +// CHECK: @local_memory_alignment_global.lds_v3i16 = internal addrspace(3) global [4 x <3 x i16>] undef, align 8 +// CHECK: @local_memory_alignment_global.lds_v4i16 = internal addrspace(3) global [4 x <4 x i16>] undef, align 8 +// CHECK: @local_memory_alignment_global.lds_v8i16 = internal addrspace(3) global [4 x <8 x i16>] undef, align 16 +// CHECK: @local_memory_alignment_global.lds_v16i16 = internal addrspace(3) global [4 x <16 x i16>] undef, align 32 +// CHECK: @local_memory_alignment_global.lds_i32 = internal addrspace(3) global [4 x i32] undef, align 4 +// CHECK: @local_memory_alignment_global.lds_v2i32 = internal addrspace(3) global [4 x <2 x i32>] undef, align 8 +// CHECK: @local_memory_alignment_global.lds_v3i32 = internal addrspace(3) global [4 x <3 x i32>] undef, align 16 +// CHECK: @local_memory_alignment_global.lds_v4i32 = internal addrspace(3) global [4 x <4 x i32>] undef, align 16 +// CHECK: @local_memory_alignment_global.lds_v8i32 = internal addrspace(3) global [4 x <8 x i32>] undef, align 32 +// CHECK: @local_memory_alignment_global.lds_v16i32 = internal addrspace(3) global [4 x <16 x i32>] undef, align 64 +// CHECK: @local_memory_alignment_global.lds_i64 = internal addrspace(3) global [4 x i64] undef, align 8 +// CHECK: @local_memory_alignment_global.lds_v2i64 = internal addrspace(3) global [4 x <2 x i64>] undef, align 16 +// CHECK: @local_memory_alignment_global.lds_v3i64 = internal addrspace(3) global [4 x <3 x i64>] undef, align 32 +// CHECK: @local_memory_alignment_global.lds_v4i64 = internal addrspace(3) global [4 x <4 x i64>] undef, align 32 +// CHECK: @local_memory_alignment_global.lds_v8i64 = internal addrspace(3) global [4 x <8 x i64>] undef, align 64 +// CHECK: @local_memory_alignment_global.lds_v16i64 = internal addrspace(3) global [4 x <16 x i64>] undef, align 128 +// CHECK: @local_memory_alignment_global.lds_f16 = internal addrspace(3) global [4 x half] undef, align 2 +// CHECK: @local_memory_alignment_global.lds_v2f16 = internal addrspace(3) global [4 x <2 x half>] undef, align 4 +// CHECK: @local_memory_alignment_global.lds_v3f16 = internal addrspace(3) global [4 x <3 x half>] undef, align 8 +// CHECK: @local_memory_alignment_global.lds_v4f16 = internal addrspace(3) global [4 x <4 x half>] undef, align 8 +// CHECK: @local_memory_alignment_global.lds_v8f16 = internal addrspace(3) global [4 x <8 x half>] undef, align 16 +// CHECK: @local_memory_alignment_global.lds_v16f16 = internal addrspace(3) global [4 x <16 x half>] undef, align 32 +// CHECK: @local_memory_alignment_global.lds_f32 = internal addrspace(3) global [4 x float] undef, align 4 +// CHECK: @local_memory_alignment_global.lds_v2f32 = internal addrspace(3) global [4 x <2 x float>] undef, align 8 +// CHECK: @local_memory_alignment_global.lds_v3f32 = internal addrspace(3) global [4 x <3 x float>] undef, align 16 +// CHECK: @local_memory_alignment_global.lds_v4f32 = internal addrspace(3) global [4 x <4 x float>] undef, align 16 +// CHECK: @local_memory_alignment_global.lds_v8f32 = internal addrspace(3) global [4 x <8 x float>] undef, align 32 +// CHECK: @local_memory_alignment_global.lds_v16f32 = internal addrspace(3) global [4 x <16 x float>] undef, align 64 +// CHECK: @local_memory_alignment_global.lds_f64 = internal addrspace(3) global [4 x double] undef, align 8 +// CHECK: @local_memory_alignment_global.lds_v2f64 = internal addrspace(3) global [4 x <2 x double>] undef, align 16 +// CHECK: @local_memory_alignment_global.lds_v3f64 = internal addrspace(3) global [4 x <3 x double>] undef, align 32 +// CHECK: @local_memory_alignment_global.lds_v4f64 = internal addrspace(3) global [4 x <4 x double>] undef, align 32 +// CHECK: @local_memory_alignment_global.lds_v8f64 = internal addrspace(3) global [4 x <8 x double>] undef, align 64 +// CHECK: @local_memory_alignment_global.lds_v16f64 = internal addrspace(3) global [4 x <16 x double>] undef, align 128 + + +// CHECK-LABEL: @local_memory_alignment_global( +// CHECK: store volatile i8 0, i8 addrspace(3)* getelementptr inbounds ([4 x i8], [4 x i8] addrspace(3)* @local_memory_alignment_global.lds_i8, i32 0, i32 0), align 1 +// CHECK: store volatile <2 x i8> zeroinitializer, <2 x i8> addrspace(3)* getelementptr inbounds ([4 x <2 x i8>], [4 x <2 x i8>] addrspace(3)* @local_memory_alignment_global.lds_v2i8, i32 0, i32 0), align 2 +// CHECK: store volatile <4 x i8> <i8 0, i8 0, i8 0, i8 undef>, <4 x i8> addrspace(3)* bitcast ([4 x <3 x i8>] addrspace(3)* @local_memory_alignment_global.lds_v3i8 to <4 x i8> addrspace(3)*), align 4 +// CHECK: store volatile <4 x i8> zeroinitializer, <4 x i8> addrspace(3)* getelementptr inbounds ([4 x <4 x i8>], [4 x <4 x i8>] addrspace(3)* @local_memory_alignment_global.lds_v4i8, i32 0, i32 0), align 4 +// CHECK: store volatile <8 x i8> zeroinitializer, <8 x i8> addrspace(3)* getelementptr inbounds ([4 x <8 x i8>], [4 x <8 x i8>] addrspace(3)* @local_memory_alignment_global.lds_v8i8, i32 0, i32 0), align 8 +// CHECK: store volatile <16 x i8> zeroinitializer, <16 x i8> addrspace(3)* getelementptr inbounds ([4 x <16 x i8>], [4 x <16 x i8>] addrspace(3)* @local_memory_alignment_global.lds_v16i8, i32 0, i32 0), align 16 +// CHECK: store volatile i16 0, i16 addrspace(3)* getelementptr inbounds ([4 x i16], [4 x i16] addrspace(3)* @local_memory_alignment_global.lds_i16, i32 0, i32 0), align 2 +// CHECK: store volatile <2 x i16> zeroinitializer, <2 x i16> addrspace(3)* getelementptr inbounds ([4 x <2 x i16>], [4 x <2 x i16>] addrspace(3)* @local_memory_alignment_global.lds_v2i16, i32 0, i32 0), align 4 +// CHECK: store volatile <4 x i16> <i16 0, i16 0, i16 0, i16 undef>, <4 x i16> addrspace(3)* bitcast ([4 x <3 x i16>] addrspace(3)* @local_memory_alignment_global.lds_v3i16 to <4 x i16> addrspace(3)*), align 8 +// CHECK: store volatile <4 x i16> zeroinitializer, <4 x i16> addrspace(3)* getelementptr inbounds ([4 x <4 x i16>], [4 x <4 x i16>] addrspace(3)* @local_memory_alignment_global.lds_v4i16, i32 0, i32 0), align 8 +// CHECK: store volatile <8 x i16> zeroinitializer, <8 x i16> addrspace(3)* getelementptr inbounds ([4 x <8 x i16>], [4 x <8 x i16>] addrspace(3)* @local_memory_alignment_global.lds_v8i16, i32 0, i32 0), align 16 +// CHECK: store volatile <16 x i16> zeroinitializer, <16 x i16> addrspace(3)* getelementptr inbounds ([4 x <16 x i16>], [4 x <16 x i16>] addrspace(3)* @local_memory_alignment_global.lds_v16i16, i32 0, i32 0), align 32 +// CHECK: store volatile i32 0, i32 addrspace(3)* getelementptr inbounds ([4 x i32], [4 x i32] addrspace(3)* @local_memory_alignment_global.lds_i32, i32 0, i32 0), align 4 +// CHECK: store volatile <2 x i32> zeroinitializer, <2 x i32> addrspace(3)* getelementptr inbounds ([4 x <2 x i32>], [4 x <2 x i32>] addrspace(3)* @local_memory_alignment_global.lds_v2i32, i32 0, i32 0), align 8 +// CHECK: store volatile <4 x i32> <i32 0, i32 0, i32 0, i32 undef>, <4 x i32> addrspace(3)* bitcast ([4 x <3 x i32>] addrspace(3)* @local_memory_alignment_global.lds_v3i32 to <4 x i32> addrspace(3)*), align 16 +// CHECK: store volatile <4 x i32> zeroinitializer, <4 x i32> addrspace(3)* getelementptr inbounds ([4 x <4 x i32>], [4 x <4 x i32>] addrspace(3)* @local_memory_alignment_global.lds_v4i32, i32 0, i32 0), align 16 +// CHECK: store volatile <8 x i32> zeroinitializer, <8 x i32> addrspace(3)* getelementptr inbounds ([4 x <8 x i32>], [4 x <8 x i32>] addrspace(3)* @local_memory_alignment_global.lds_v8i32, i32 0, i32 0), align 32 +// CHECK: store volatile <16 x i32> zeroinitializer, <16 x i32> addrspace(3)* getelementptr inbounds ([4 x <16 x i32>], [4 x <16 x i32>] addrspace(3)* @local_memory_alignment_global.lds_v16i32, i32 0, i32 0), align 64 +// CHECK: store volatile i64 0, i64 addrspace(3)* getelementptr inbounds ([4 x i64], [4 x i64] addrspace(3)* @local_memory_alignment_global.lds_i64, i32 0, i32 0), align 8 +// CHECK: store volatile <2 x i64> zeroinitializer, <2 x i64> addrspace(3)* getelementptr inbounds ([4 x <2 x i64>], [4 x <2 x i64>] addrspace(3)* @local_memory_alignment_global.lds_v2i64, i32 0, i32 0), align 16 +// CHECK: store volatile <4 x i64> <i64 0, i64 0, i64 0, i64 undef>, <4 x i64> addrspace(3)* bitcast ([4 x <3 x i64>] addrspace(3)* @local_memory_alignment_global.lds_v3i64 to <4 x i64> addrspace(3)*), align 32 +// CHECK: store volatile <4 x i64> zeroinitializer, <4 x i64> addrspace(3)* getelementptr inbounds ([4 x <4 x i64>], [4 x <4 x i64>] addrspace(3)* @local_memory_alignment_global.lds_v4i64, i32 0, i32 0), align 32 +// CHECK: store volatile <8 x i64> zeroinitializer, <8 x i64> addrspace(3)* getelementptr inbounds ([4 x <8 x i64>], [4 x <8 x i64>] addrspace(3)* @local_memory_alignment_global.lds_v8i64, i32 0, i32 0), align 64 +// CHECK: store volatile <16 x i64> zeroinitializer, <16 x i64> addrspace(3)* getelementptr inbounds ([4 x <16 x i64>], [4 x <16 x i64>] addrspace(3)* @local_memory_alignment_global.lds_v16i64, i32 0, i32 0), align 128 +// CHECK: store volatile half 0xH0000, half addrspace(3)* getelementptr inbounds ([4 x half], [4 x half] addrspace(3)* @local_memory_alignment_global.lds_f16, i32 0, i32 0), align 2 +// CHECK: store volatile <2 x half> zeroinitializer, <2 x half> addrspace(3)* getelementptr inbounds ([4 x <2 x half>], [4 x <2 x half>] addrspace(3)* @local_memory_alignment_global.lds_v2f16, i32 0, i32 0), align 4 +// CHECK: store volatile <4 x half> <half 0xH0000, half 0xH0000, half 0xH0000, half undef>, <4 x half> addrspace(3)* bitcast ([4 x <3 x half>] addrspace(3)* @local_memory_alignment_global.lds_v3f16 to <4 x half> addrspace(3)*), align 8 +// CHECK: store volatile <4 x half> zeroinitializer, <4 x half> addrspace(3)* getelementptr inbounds ([4 x <4 x half>], [4 x <4 x half>] addrspace(3)* @local_memory_alignment_global.lds_v4f16, i32 0, i32 0), align 8 +// CHECK: store volatile <8 x half> zeroinitializer, <8 x half> addrspace(3)* getelementptr inbounds ([4 x <8 x half>], [4 x <8 x half>] addrspace(3)* @local_memory_alignment_global.lds_v8f16, i32 0, i32 0), align 16 +// CHECK: store volatile <16 x half> zeroinitializer, <16 x half> addrspace(3)* getelementptr inbounds ([4 x <16 x half>], [4 x <16 x half>] addrspace(3)* @local_memory_alignment_global.lds_v16f16, i32 0, i32 0), align 32 +// CHECK: store volatile float 0.000000e+00, float addrspace(3)* getelementptr inbounds ([4 x float], [4 x float] addrspace(3)* @local_memory_alignment_global.lds_f32, i32 0, i32 0), align 4 +// CHECK: store volatile <2 x float> zeroinitializer, <2 x float> addrspace(3)* getelementptr inbounds ([4 x <2 x float>], [4 x <2 x float>] addrspace(3)* @local_memory_alignment_global.lds_v2f32, i32 0, i32 0), align 8 +// CHECK: store volatile <4 x float> <float 0.000000e+00, float 0.000000e+00, float 0.000000e+00, float undef>, <4 x float> addrspace(3)* bitcast ([4 x <3 x float>] addrspace(3)* @local_memory_alignment_global.lds_v3f32 to <4 x float> addrspace(3)*), align 16 +// CHECK: store volatile <4 x float> zeroinitializer, <4 x float> addrspace(3)* getelementptr inbounds ([4 x <4 x float>], [4 x <4 x float>] addrspace(3)* @local_memory_alignment_global.lds_v4f32, i32 0, i32 0), align 16 +// CHECK: store volatile <8 x float> zeroinitializer, <8 x float> addrspace(3)* getelementptr inbounds ([4 x <8 x float>], [4 x <8 x float>] addrspace(3)* @local_memory_alignment_global.lds_v8f32, i32 0, i32 0), align 32 +// CHECK: store volatile <16 x float> zeroinitializer, <16 x float> addrspace(3)* getelementptr inbounds ([4 x <16 x float>], [4 x <16 x float>] addrspace(3)* @local_memory_alignment_global.lds_v16f32, i32 0, i32 0), align 64 +// CHECK: store volatile double 0.000000e+00, double addrspace(3)* getelementptr inbounds ([4 x double], [4 x double] addrspace(3)* @local_memory_alignment_global.lds_f64, i32 0, i32 0), align 8 +// CHECK: store volatile <2 x double> zeroinitializer, <2 x double> addrspace(3)* getelementptr inbounds ([4 x <2 x double>], [4 x <2 x double>] addrspace(3)* @local_memory_alignment_global.lds_v2f64, i32 0, i32 0), align 16 +// CHECK: store volatile <4 x double> <double 0.000000e+00, double 0.000000e+00, double 0.000000e+00, double undef>, <4 x double> addrspace(3)* bitcast ([4 x <3 x double>] addrspace(3)* @local_memory_alignment_global.lds_v3f64 to <4 x double> addrspace(3)*), align 32 +// CHECK: store volatile <4 x double> zeroinitializer, <4 x double> addrspace(3)* getelementptr inbounds ([4 x <4 x double>], [4 x <4 x double>] addrspace(3)* @local_memory_alignment_global.lds_v4f64, i32 0, i32 0), align 32 +// CHECK: store volatile <8 x double> zeroinitializer, <8 x double> addrspace(3)* getelementptr inbounds ([4 x <8 x double>], [4 x <8 x double>] addrspace(3)* @local_memory_alignment_global.lds_v8f64, i32 0, i32 0), align 64 +// CHECK: store volatile <16 x double> zeroinitializer, <16 x double> addrspace(3)* getelementptr inbounds ([4 x <16 x double>], [4 x <16 x double>] addrspace(3)* @local_memory_alignment_global.lds_v16f64, i32 0, i32 0), align 128 +kernel void local_memory_alignment_global() +{ + volatile local char lds_i8[4]; + volatile local char2 lds_v2i8[4]; + volatile local char3 lds_v3i8[4]; + volatile local char4 lds_v4i8[4]; + volatile local char8 lds_v8i8[4]; + volatile local char16 lds_v16i8[4]; + + volatile local short lds_i16[4]; + volatile local short2 lds_v2i16[4]; + volatile local short3 lds_v3i16[4]; + volatile local short4 lds_v4i16[4]; + volatile local short8 lds_v8i16[4]; + volatile local short16 lds_v16i16[4]; + + volatile local int lds_i32[4]; + volatile local int2 lds_v2i32[4]; + volatile local int3 lds_v3i32[4]; + volatile local int4 lds_v4i32[4]; + volatile local int8 lds_v8i32[4]; + volatile local int16 lds_v16i32[4]; + + volatile local long lds_i64[4]; + volatile local long2 lds_v2i64[4]; + volatile local long3 lds_v3i64[4]; + volatile local long4 lds_v4i64[4]; + volatile local long8 lds_v8i64[4]; + volatile local long16 lds_v16i64[4]; + + volatile local half lds_f16[4]; + volatile local half2 lds_v2f16[4]; + volatile local half3 lds_v3f16[4]; + volatile local half4 lds_v4f16[4]; + volatile local half8 lds_v8f16[4]; + volatile local half16 lds_v16f16[4]; + + volatile local float lds_f32[4]; + volatile local float2 lds_v2f32[4]; + volatile local float3 lds_v3f32[4]; + volatile local float4 lds_v4f32[4]; + volatile local float8 lds_v8f32[4]; + volatile local float16 lds_v16f32[4]; + + volatile local double lds_f64[4]; + volatile local double2 lds_v2f64[4]; + volatile local double3 lds_v3f64[4]; + volatile local double4 lds_v4f64[4]; + volatile local double8 lds_v8f64[4]; + volatile local double16 lds_v16f64[4]; + + *lds_i8 = 0; + *lds_v2i8 = 0; + *lds_v3i8 = 0; + *lds_v4i8 = 0; + *lds_v8i8 = 0; + *lds_v16i8 = 0; + + *lds_i16 = 0; + *lds_v2i16 = 0; + *lds_v3i16 = 0; + *lds_v4i16 = 0; + *lds_v8i16 = 0; + *lds_v16i16 = 0; + + *lds_i32 = 0; + *lds_v2i32 = 0; + *lds_v3i32 = 0; + *lds_v4i32 = 0; + *lds_v8i32 = 0; + *lds_v16i32 = 0; + + *lds_i64 = 0; + *lds_v2i64 = 0; + *lds_v3i64 = 0; + *lds_v4i64 = 0; + *lds_v8i64 = 0; + *lds_v16i64 = 0; + + *lds_f16 = 0; + *lds_v2f16 = 0; + *lds_v3f16 = 0; + *lds_v4f16 = 0; + *lds_v8f16 = 0; + *lds_v16f16 = 0; + + *lds_f32 = 0; + *lds_v2f32 = 0; + *lds_v3f32 = 0; + *lds_v4f32 = 0; + *lds_v8f32 = 0; + *lds_v16f32 = 0; + + *lds_f64 = 0; + *lds_v2f64 = 0; + *lds_v3f64 = 0; + *lds_v4f64 = 0; + *lds_v8f64 = 0; + *lds_v16f64 = 0; +} + +kernel void local_memory_alignment_arg( + volatile local char* lds_i8, + volatile local char2* lds_v2i8, + volatile local char3* lds_v3i8, + volatile local char4* lds_v4i8, + volatile local char8* lds_v8i8, + volatile local char16* lds_v16i8, + + volatile local short* lds_i16, + volatile local short2* lds_v2i16, + volatile local short3* lds_v3i16, + volatile local short4* lds_v4i16, + volatile local short8* lds_v8i16, + volatile local short16* lds_v16i16, + + volatile local int* lds_i32, + volatile local int2* lds_v2i32, + volatile local int3* lds_v3i32, + volatile local int4* lds_v4i32, + volatile local int8* lds_v8i32, + volatile local int16* lds_v16i32, + + volatile local long* lds_i64, + volatile local long2* lds_v2i64, + volatile local long3* lds_v3i64, + volatile local long4* lds_v4i64, + volatile local long8* lds_v8i64, + volatile local long16* lds_v16i64, + + volatile local half* lds_f16, + volatile local half2* lds_v2f16, + volatile local half3* lds_v3f16, + volatile local half4* lds_v4f16, + volatile local half8* lds_v8f16, + volatile local half16* lds_v16f16, + + volatile local float* lds_f32, + volatile local float2* lds_v2f32, + volatile local float3* lds_v3f32, + volatile local float4* lds_v4f32, + volatile local float8* lds_v8f32, + volatile local float16* lds_v16f32, + + volatile local double* lds_f64, + volatile local double2* lds_v2f64, + volatile local double3* lds_v3f64, + volatile local double4* lds_v4f64, + volatile local double8* lds_v8f64, + volatile local double16* lds_v16f64) +{ + *lds_i8 = 0; + *lds_v2i8 = 0; + *lds_v3i8 = 0; + *lds_v4i8 = 0; + *lds_v8i8 = 0; + *lds_v16i8 = 0; + + *lds_i16 = 0; + *lds_v2i16 = 0; + *lds_v3i16 = 0; + *lds_v4i16 = 0; + *lds_v8i16 = 0; + *lds_v16i16 = 0; + + *lds_i32 = 0; + *lds_v2i32 = 0; + *lds_v3i32 = 0; + *lds_v4i32 = 0; + *lds_v8i32 = 0; + *lds_v16i32 = 0; + + *lds_i64 = 0; + *lds_v2i64 = 0; + *lds_v3i64 = 0; + *lds_v4i64 = 0; + *lds_v8i64 = 0; + *lds_v16i64 = 0; + + *lds_f16 = 0; + *lds_v2f16 = 0; + *lds_v3f16 = 0; + *lds_v4f16 = 0; + *lds_v8f16 = 0; + *lds_v16f16 = 0; + + *lds_f32 = 0; + *lds_v2f32 = 0; + *lds_v3f32 = 0; + *lds_v4f32 = 0; + *lds_v8f32 = 0; + *lds_v16f32 = 0; + + *lds_f64 = 0; + *lds_v2f64 = 0; + *lds_v3f64 = 0; + *lds_v4f64 = 0; + *lds_v8f64 = 0; + *lds_v16f64 = 0; +} + +// CHECK-LABEL: @private_memory_alignment_alloca( +// CHECK: %private_i8 = alloca [4 x i8], align 1 +// CHECK: %private_v2i8 = alloca [4 x <2 x i8>], align 2 +// CHECK: %private_v3i8 = alloca [4 x <3 x i8>], align 4 +// CHECK: %private_v4i8 = alloca [4 x <4 x i8>], align 4 +// CHECK: %private_v8i8 = alloca [4 x <8 x i8>], align 8 +// CHECK: %private_v16i8 = alloca [4 x <16 x i8>], align 16 +// CHECK: %private_i16 = alloca [4 x i16], align 2 +// CHECK: %private_v2i16 = alloca [4 x <2 x i16>], align 4 +// CHECK: %private_v3i16 = alloca [4 x <3 x i16>], align 8 +// CHECK: %private_v4i16 = alloca [4 x <4 x i16>], align 8 +// CHECK: %private_v8i16 = alloca [4 x <8 x i16>], align 16 +// CHECK: %private_v16i16 = alloca [4 x <16 x i16>], align 32 +// CHECK: %private_i32 = alloca [4 x i32], align 4 +// CHECK: %private_v2i32 = alloca [4 x <2 x i32>], align 8 +// CHECK: %private_v3i32 = alloca [4 x <3 x i32>], align 16 +// CHECK: %private_v4i32 = alloca [4 x <4 x i32>], align 16 +// CHECK: %private_v8i32 = alloca [4 x <8 x i32>], align 32 +// CHECK: %private_v16i32 = alloca [4 x <16 x i32>], align 64 +// CHECK: %private_i64 = alloca [4 x i64], align 8 +// CHECK: %private_v2i64 = alloca [4 x <2 x i64>], align 16 +// CHECK: %private_v3i64 = alloca [4 x <3 x i64>], align 32 +// CHECK: %private_v4i64 = alloca [4 x <4 x i64>], align 32 +// CHECK: %private_v8i64 = alloca [4 x <8 x i64>], align 64 +// CHECK: %private_v16i64 = alloca [4 x <16 x i64>], align 128 +// CHECK: %private_f16 = alloca [4 x half], align 2 +// CHECK: %private_v2f16 = alloca [4 x <2 x half>], align 4 +// CHECK: %private_v3f16 = alloca [4 x <3 x half>], align 8 +// CHECK: %private_v4f16 = alloca [4 x <4 x half>], align 8 +// CHECK: %private_v8f16 = alloca [4 x <8 x half>], align 16 +// CHECK: %private_v16f16 = alloca [4 x <16 x half>], align 32 +// CHECK: %private_f32 = alloca [4 x float], align 4 +// CHECK: %private_v2f32 = alloca [4 x <2 x float>], align 8 +// CHECK: %private_v3f32 = alloca [4 x <3 x float>], align 16 +// CHECK: %private_v4f32 = alloca [4 x <4 x float>], align 16 +// CHECK: %private_v8f32 = alloca [4 x <8 x float>], align 32 +// CHECK: %private_v16f32 = alloca [4 x <16 x float>], align 64 +// CHECK: %private_f64 = alloca [4 x double], align 8 +// CHECK: %private_v2f64 = alloca [4 x <2 x double>], align 16 +// CHECK: %private_v3f64 = alloca [4 x <3 x double>], align 32 +// CHECK: %private_v4f64 = alloca [4 x <4 x double>], align 32 +// CHECK: %private_v8f64 = alloca [4 x <8 x double>], align 64 +// CHECK: %private_v16f64 = alloca [4 x <16 x double>], align 128 + +// CHECK: store volatile i8 0, i8* %arraydecay, align 1 +// CHECK: store volatile <2 x i8> zeroinitializer, <2 x i8>* %arraydecay{{[0-9]+}}, align 2 +// CHECK: store volatile <4 x i8> <i8 0, i8 0, i8 0, i8 undef>, <4 x i8>* %storetmp, align 4 +// CHECK: store volatile <4 x i8> zeroinitializer, <4 x i8>* %arraydecay{{[0-9]+}}, align 4 +// CHECK: store volatile <8 x i8> zeroinitializer, <8 x i8>* %arraydecay{{[0-9]+}}, align 8 +// CHECK: store volatile <16 x i8> zeroinitializer, <16 x i8>* %arraydecay{{[0-9]+}}, align 16 +// CHECK: store volatile i16 0, i16* %arraydecay{{[0-9]+}}, align 2 +// CHECK: store volatile <2 x i16> zeroinitializer, <2 x i16>* %arraydecay{{[0-9]+}}, align 4 +// CHECK: store volatile <4 x i16> <i16 0, i16 0, i16 0, i16 undef>, <4 x i16>* %storetmp{{[0-9]+}}, align 8 +// CHECK: store volatile <4 x i16> zeroinitializer, <4 x i16>* %arraydecay{{[0-9]+}}, align 8 +// CHECK: store volatile <8 x i16> zeroinitializer, <8 x i16>* %arraydecay{{[0-9]+}}, align 16 +// CHECK: store volatile <16 x i16> zeroinitializer, <16 x i16>* %arraydecay{{[0-9]+}}, align 32 +// CHECK: store volatile i32 0, i32* %arraydecay{{[0-9]+}}, align 4 +// CHECK: store volatile <2 x i32> zeroinitializer, <2 x i32>* %arraydecay{{[0-9]+}}, align 8 +// CHECK: store volatile <4 x i32> <i32 0, i32 0, i32 0, i32 undef>, <4 x i32>* %storetmp16, align 16 +// CHECK: store volatile <4 x i32> zeroinitializer, <4 x i32>* %arraydecay{{[0-9]+}}, align 16 +// CHECK: store volatile <8 x i32> zeroinitializer, <8 x i32>* %arraydecay{{[0-9]+}}, align 32 +// CHECK: store volatile <16 x i32> zeroinitializer, <16 x i32>* %arraydecay{{[0-9]+}}, align 64 +// CHECK: store volatile i64 0, i64* %arraydecay{{[0-9]+}}, align 8 +// CHECK: store volatile <2 x i64> zeroinitializer, <2 x i64>* %arraydecay{{[0-9]+}}, align 16 +// CHECK: store volatile <4 x i64> <i64 0, i64 0, i64 0, i64 undef>, <4 x i64>* %storetmp23, align 32 +// CHECK: store volatile <4 x i64> zeroinitializer, <4 x i64>* %arraydecay{{[0-9]+}}, align 32 +// CHECK: store volatile <8 x i64> zeroinitializer, <8 x i64>* %arraydecay{{[0-9]+}}, align 64 +// CHECK: store volatile <16 x i64> zeroinitializer, <16 x i64>* %arraydecay{{[0-9]+}}, align 128 +// CHECK: store volatile half 0xH0000, half* %arraydecay{{[0-9]+}}, align 2 +// CHECK: store volatile <2 x half> zeroinitializer, <2 x half>* %arraydecay{{[0-9]+}}, align 4 +// CHECK: store volatile <4 x half> <half 0xH0000, half 0xH0000, half 0xH0000, half undef>, <4 x half>* %storetmp{{[0-9]+}}, align 8 +// CHECK: store volatile <4 x half> zeroinitializer, <4 x half>* %arraydecay{{[0-9]+}}, align 8 +// CHECK: store volatile <8 x half> zeroinitializer, <8 x half>* %arraydecay{{[0-9]+}}, align 16 +// CHECK: store volatile <16 x half> zeroinitializer, <16 x half>* %arraydecay{{[0-9]+}}, align 32 +// CHECK: store volatile float 0.000000e+00, float* %arraydecay34, align 4 +// CHECK: store volatile <2 x float> zeroinitializer, <2 x float>* %arraydecay{{[0-9]+}}, align 8 +// CHECK: store volatile <4 x float> <float 0.000000e+00, float 0.000000e+00, float 0.000000e+00, float undef>, <4 x float>* %storetmp{{[0-9]+}}, align 16 +// CHECK: store volatile <4 x float> zeroinitializer, <4 x float>* %arraydecay{{[0-9]+}}, align 16 +// CHECK: store volatile <8 x float> zeroinitializer, <8 x float>* %arraydecay{{[0-9]+}}, align 32 +// CHECK: store volatile <16 x float> zeroinitializer, <16 x float>* %arraydecay{{[0-9]+}}, align 64 +// CHECK: store volatile double 0.000000e+00, double* %arraydecay{{[0-9]+}}, align 8 +// CHECK: store volatile <2 x double> zeroinitializer, <2 x double>* %arraydecay{{[0-9]+}}, align 16 +// CHECK: store volatile <4 x double> <double 0.000000e+00, double 0.000000e+00, double 0.000000e+00, double undef>, <4 x double>* %storetmp{{[0-9]+}}, align 32 +// CHECK: store volatile <4 x double> zeroinitializer, <4 x double>* %arraydecay{{[0-9]+}}, align 32 +// CHECK: store volatile <8 x double> zeroinitializer, <8 x double>* %arraydecay{{[0-9]+}}, align 64 +// CHECK: store volatile <16 x double> zeroinitializer, <16 x double>* %arraydecay{{[0-9]+}}, align 128 +kernel void private_memory_alignment_alloca() +{ + volatile private char private_i8[4]; + volatile private char2 private_v2i8[4]; + volatile private char3 private_v3i8[4]; + volatile private char4 private_v4i8[4]; + volatile private char8 private_v8i8[4]; + volatile private char16 private_v16i8[4]; + + volatile private short private_i16[4]; + volatile private short2 private_v2i16[4]; + volatile private short3 private_v3i16[4]; + volatile private short4 private_v4i16[4]; + volatile private short8 private_v8i16[4]; + volatile private short16 private_v16i16[4]; + + volatile private int private_i32[4]; + volatile private int2 private_v2i32[4]; + volatile private int3 private_v3i32[4]; + volatile private int4 private_v4i32[4]; + volatile private int8 private_v8i32[4]; + volatile private int16 private_v16i32[4]; + + volatile private long private_i64[4]; + volatile private long2 private_v2i64[4]; + volatile private long3 private_v3i64[4]; + volatile private long4 private_v4i64[4]; + volatile private long8 private_v8i64[4]; + volatile private long16 private_v16i64[4]; + + volatile private half private_f16[4]; + volatile private half2 private_v2f16[4]; + volatile private half3 private_v3f16[4]; + volatile private half4 private_v4f16[4]; + volatile private half8 private_v8f16[4]; + volatile private half16 private_v16f16[4]; + + volatile private float private_f32[4]; + volatile private float2 private_v2f32[4]; + volatile private float3 private_v3f32[4]; + volatile private float4 private_v4f32[4]; + volatile private float8 private_v8f32[4]; + volatile private float16 private_v16f32[4]; + + volatile private double private_f64[4]; + volatile private double2 private_v2f64[4]; + volatile private double3 private_v3f64[4]; + volatile private double4 private_v4f64[4]; + volatile private double8 private_v8f64[4]; + volatile private double16 private_v16f64[4]; + + *private_i8 = 0; + *private_v2i8 = 0; + *private_v3i8 = 0; + *private_v4i8 = 0; + *private_v8i8 = 0; + *private_v16i8 = 0; + + *private_i16 = 0; + *private_v2i16 = 0; + *private_v3i16 = 0; + *private_v4i16 = 0; + *private_v8i16 = 0; + *private_v16i16 = 0; + + *private_i32 = 0; + *private_v2i32 = 0; + *private_v3i32 = 0; + *private_v4i32 = 0; + *private_v8i32 = 0; + *private_v16i32 = 0; + + *private_i64 = 0; + *private_v2i64 = 0; + *private_v3i64 = 0; + *private_v4i64 = 0; + *private_v8i64 = 0; + *private_v16i64 = 0; + + *private_f16 = 0; + *private_v2f16 = 0; + *private_v3f16 = 0; + *private_v4f16 = 0; + *private_v8f16 = 0; + *private_v16f16 = 0; + + *private_f32 = 0; + *private_v2f32 = 0; + *private_v3f32 = 0; + *private_v4f32 = 0; + *private_v8f32 = 0; + *private_v16f32 = 0; + + *private_f64 = 0; + *private_v2f64 = 0; + *private_v3f64 = 0; + *private_v4f64 = 0; + *private_v8f64 = 0; + *private_v16f64 = 0; +} diff --git a/test/CodeGenOpenCL/amdgpu-attrs.cl b/test/CodeGenOpenCL/amdgpu-attrs.cl index 4ca85d316104..c914f2e6514f 100644 --- a/test/CodeGenOpenCL/amdgpu-attrs.cl +++ b/test/CodeGenOpenCL/amdgpu-attrs.cl @@ -129,6 +129,16 @@ 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]+]] } +__attribute__((reqd_work_group_size(32, 2, 1))) // expected-no-diagnostics +kernel void reqd_work_group_size_32_2_1() { +// CHECK: define amdgpu_kernel void @reqd_work_group_size_32_2_1() [[FLAT_WORK_GROUP_SIZE_64_64:#[0-9]+]] +} +__attribute__((reqd_work_group_size(32, 2, 1), amdgpu_flat_work_group_size(16, 128))) // expected-no-diagnostics +kernel void reqd_work_group_size_32_2_1_flat_work_group_size_16_128() { +// CHECK: define amdgpu_kernel void @reqd_work_group_size_32_2_1_flat_work_group_size_16_128() [[FLAT_WORK_GROUP_SIZE_16_128:#[0-9]+]] +} + + // Make sure this is silently accepted on other targets. // X86-NOT: "amdgpu-flat-work-group-size" // X86-NOT: "amdgpu-waves-per-eu" @@ -142,6 +152,8 @@ kernel void flat_work_group_size_32_64_waves_per_eu_2_4_num_sgpr_32_num_vgpr_64( // 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 [[FLAT_WORK_GROUP_SIZE_64_64]] = { noinline nounwind "amdgpu-flat-work-group-size"="64,64" +// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_16_128]] = { noinline nounwind "amdgpu-flat-work-group-size"="16,128" // 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" diff --git a/test/CodeGenOpenCL/amdgpu-debug-info-pointer-address-space.cl b/test/CodeGenOpenCL/amdgpu-debug-info-pointer-address-space.cl new file mode 100644 index 000000000000..c0b30095a679 --- /dev/null +++ b/test/CodeGenOpenCL/amdgpu-debug-info-pointer-address-space.cl @@ -0,0 +1,125 @@ +// RUN: %clang -cl-std=CL2.0 -emit-llvm -g -O0 -S -target amdgcn-amd-amdhsa -mcpu=fiji -o - %s | FileCheck %s + +// CHECK-DAG: ![[DWARF_ADDRESS_SPACE_NONE:[0-9]+]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !{{[0-9]+}}, size: {{[0-9]+}}) +// CHECK-DAG: ![[DWARF_ADDRESS_SPACE_LOCAL:[0-9]+]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !{{[0-9]+}}, size: {{[0-9]+}}, dwarfAddressSpace: 2) +// CHECK-DAG: ![[DWARF_ADDRESS_SPACE_PRIVATE:[0-9]+]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !{{[0-9]+}}, size: {{[0-9]+}}, dwarfAddressSpace: 1) + +// CHECK-DAG: distinct !DIGlobalVariable(name: "FileVar0", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]], isLocal: false, isDefinition: true) +global int *FileVar0; +// CHECK-DAG: distinct !DIGlobalVariable(name: "FileVar1", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]], isLocal: false, isDefinition: true) +constant int *FileVar1; +// CHECK-DAG: distinct !DIGlobalVariable(name: "FileVar2", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_LOCAL]], isLocal: false, isDefinition: true) +local int *FileVar2; +// CHECK-DAG: distinct !DIGlobalVariable(name: "FileVar3", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_PRIVATE]], isLocal: false, isDefinition: true) +private int *FileVar3; +// CHECK-DAG: distinct !DIGlobalVariable(name: "FileVar4", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]], isLocal: false, isDefinition: true) +int *FileVar4; + +// CHECK-DAG: distinct !DIGlobalVariable(name: "FileVar5", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]], isLocal: false, isDefinition: true) +global int *global FileVar5; +// CHECK-DAG: distinct !DIGlobalVariable(name: "FileVar6", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]], isLocal: false, isDefinition: true) +constant int *global FileVar6; +// CHECK-DAG: distinct !DIGlobalVariable(name: "FileVar7", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_LOCAL]], isLocal: false, isDefinition: true) +local int *global FileVar7; +// CHECK-DAG: distinct !DIGlobalVariable(name: "FileVar8", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_PRIVATE]], isLocal: false, isDefinition: true) +private int *global FileVar8; +// CHECK-DAG: distinct !DIGlobalVariable(name: "FileVar9", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]], isLocal: false, isDefinition: true) +int *global FileVar9; + +// CHECK-DAG: distinct !DIGlobalVariable(name: "FileVar10", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]], isLocal: false, isDefinition: true) +global int *constant FileVar10 = 0; +// CHECK-DAG: distinct !DIGlobalVariable(name: "FileVar11", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]], isLocal: false, isDefinition: true) +constant int *constant FileVar11 = 0; +// CHECK-DAG: distinct !DIGlobalVariable(name: "FileVar12", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_LOCAL]], isLocal: false, isDefinition: true) +local int *constant FileVar12 = 0; +// CHECK-DAG: distinct !DIGlobalVariable(name: "FileVar13", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_PRIVATE]], isLocal: false, isDefinition: true) +private int *constant FileVar13 = 0; +// CHECK-DAG: distinct !DIGlobalVariable(name: "FileVar14", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]], isLocal: false, isDefinition: true) +int *constant FileVar14 = 0; + +kernel void kernel1( + // CHECK-DAG: !DILocalVariable(name: "KernelArg0", arg: {{[0-9]+}}, scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]]) + global int *KernelArg0, + // CHECK-DAG: !DILocalVariable(name: "KernelArg1", arg: {{[0-9]+}}, scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]]) + constant int *KernelArg1, + // CHECK-DAG: !DILocalVariable(name: "KernelArg2", arg: {{[0-9]+}}, scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_LOCAL]]) + local int *KernelArg2) { + private int *Tmp0; + int *Tmp1; + + // CHECK-DAG: !DILocalVariable(name: "FuncVar0", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]]) + global int *FuncVar0 = KernelArg0; + // CHECK-DAG: !DILocalVariable(name: "FuncVar1", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]]) + constant int *FuncVar1 = KernelArg1; + // CHECK-DAG: !DILocalVariable(name: "FuncVar2", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_LOCAL]]) + local int *FuncVar2 = KernelArg2; + // CHECK-DAG: !DILocalVariable(name: "FuncVar3", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_PRIVATE]]) + private int *FuncVar3 = Tmp0; + // CHECK-DAG: !DILocalVariable(name: "FuncVar4", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]]) + int *FuncVar4 = Tmp1; + + // CHECK-DAG: !DILocalVariable(name: "FuncVar5", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]]) + global int *constant FuncVar5 = KernelArg0; + // CHECK-DAG: !DILocalVariable(name: "FuncVar6", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]]) + constant int *constant FuncVar6 = KernelArg1; + // CHECK-DAG: !DILocalVariable(name: "FuncVar7", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_LOCAL]]) + local int *constant FuncVar7 = KernelArg2; + // CHECK-DAG: !DILocalVariable(name: "FuncVar8", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_PRIVATE]]) + private int *constant FuncVar8 = Tmp0; + // CHECK-DAG: !DILocalVariable(name: "FuncVar9", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]]) + int *constant FuncVar9 = Tmp1; + + // CHECK-DAG: distinct !DIGlobalVariable(name: "FuncVar10", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]], isLocal: true, isDefinition: true) + global int *local FuncVar10; FuncVar10 = KernelArg0; + // CHECK-DAG: distinct !DIGlobalVariable(name: "FuncVar11", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]], isLocal: true, isDefinition: true) + constant int *local FuncVar11; FuncVar11 = KernelArg1; + // CHECK-DAG: distinct !DIGlobalVariable(name: "FuncVar12", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_LOCAL]], isLocal: true, isDefinition: true) + local int *local FuncVar12; FuncVar12 = KernelArg2; + // CHECK-DAG: distinct !DIGlobalVariable(name: "FuncVar13", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_PRIVATE]], isLocal: true, isDefinition: true) + private int *local FuncVar13; FuncVar13 = Tmp0; + // CHECK-DAG: distinct !DIGlobalVariable(name: "FuncVar14", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]], isLocal: true, isDefinition: true) + int *local FuncVar14; FuncVar14 = Tmp1; + + // CHECK-DAG: !DILocalVariable(name: "FuncVar15", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]]) + global int *private FuncVar15 = KernelArg0; + // CHECK-DAG: !DILocalVariable(name: "FuncVar16", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]]) + constant int *private FuncVar16 = KernelArg1; + // CHECK-DAG: !DILocalVariable(name: "FuncVar17", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_LOCAL]]) + local int *private FuncVar17 = KernelArg2; + // CHECK-DAG: !DILocalVariable(name: "FuncVar18", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_PRIVATE]]) + private int *private FuncVar18 = Tmp0; + // CHECK-DAG: !DILocalVariable(name: "FuncVar19", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_ADDRESS_SPACE_NONE]]) + int *private FuncVar19 = Tmp1; +} + +struct FileStruct0 { + // CHECK-DAG: !DIDerivedType(tag: DW_TAG_member, name: "StructMem0", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, baseType: ![[DWARF_ADDRESS_SPACE_NONE]], size: {{[0-9]+}}) + global int *StructMem0; + // CHECK-DAG: !DIDerivedType(tag: DW_TAG_member, name: "StructMem1", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, baseType: ![[DWARF_ADDRESS_SPACE_NONE]], size: {{[0-9]+}}, offset: {{[0-9]+}}) + constant int *StructMem1; + // CHECK-DAG: !DIDerivedType(tag: DW_TAG_member, name: "StructMem2", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, baseType: ![[DWARF_ADDRESS_SPACE_LOCAL]], size: {{[0-9]+}}, offset: {{[0-9]+}}) + local int *StructMem2; + // CHECK-DAG: !DIDerivedType(tag: DW_TAG_member, name: "StructMem3", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, baseType: ![[DWARF_ADDRESS_SPACE_PRIVATE]], size: {{[0-9]+}}, offset: {{[0-9]+}}) + private int *StructMem3; + // CHECK-DAG: !DIDerivedType(tag: DW_TAG_member, name: "StructMem4", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, baseType: ![[DWARF_ADDRESS_SPACE_NONE]], size: {{[0-9]+}}, offset: {{[0-9]+}}) + int *StructMem4; +}; + +struct FileStruct1 { + union { + // CHECK-DAG: !DIDerivedType(tag: DW_TAG_member, name: "UnionMem0", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, baseType: ![[DWARF_ADDRESS_SPACE_NONE]], size: {{[0-9]+}}) + global int *UnionMem0; + // CHECK-DAG: !DIDerivedType(tag: DW_TAG_member, name: "UnionMem1", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, baseType: ![[DWARF_ADDRESS_SPACE_NONE]], size: {{[0-9]+}}) + constant int *UnionMem1; + // CHECK-DAG: !DIDerivedType(tag: DW_TAG_member, name: "UnionMem2", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, baseType: ![[DWARF_ADDRESS_SPACE_LOCAL]], size: {{[0-9]+}}) + local int *UnionMem2; + // CHECK-DAG: !DIDerivedType(tag: DW_TAG_member, name: "UnionMem3", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, baseType: ![[DWARF_ADDRESS_SPACE_PRIVATE]], size: {{[0-9]+}}) + private int *UnionMem3; + // CHECK-DAG: !DIDerivedType(tag: DW_TAG_member, name: "UnionMem4", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, baseType: ![[DWARF_ADDRESS_SPACE_NONE]], size: {{[0-9]+}}) + int *UnionMem4; + }; + long StructMem0; +}; + +kernel void kernel2(global struct FileStruct0 *Kernel2Arg0, + global struct FileStruct1 *Kernel2Arg1) {} diff --git a/test/CodeGenOpenCL/amdgpu-debug-info-variable-expression.cl b/test/CodeGenOpenCL/amdgpu-debug-info-variable-expression.cl new file mode 100644 index 000000000000..b02ad46224b0 --- /dev/null +++ b/test/CodeGenOpenCL/amdgpu-debug-info-variable-expression.cl @@ -0,0 +1,131 @@ +// RUN: %clang -cl-std=CL2.0 -emit-llvm -g -O0 -S -target amdgcn-amd-amdhsa -mcpu=fiji -o - %s | FileCheck %s + +// CHECK-DAG: ![[NONE:[0-9]+]] = !DIExpression() +// CHECK-DAG: ![[LOCAL:[0-9]+]] = !DIExpression(DW_OP_constu, 2, DW_OP_swap, DW_OP_xderef) +// CHECK-DAG: ![[PRIVATE:[0-9]+]] = !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef) + +// CHECK-DAG: ![[FILEVAR0:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar0", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true) +// CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR0]]) +global int *FileVar0; +// CHECK-DAG: ![[FILEVAR1:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar1", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true) +// CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR1]]) +constant int *FileVar1; +// CHECK-DAG: ![[FILEVAR2:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar2", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true) +// CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR2]]) +local int *FileVar2; +// CHECK-DAG: ![[FILEVAR3:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar3", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true) +// CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR3]]) +private int *FileVar3; +// CHECK-DAG: ![[FILEVAR4:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar4", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true) +// CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR4]]) +int *FileVar4; + +// CHECK-DAG: ![[FILEVAR5:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar5", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true) +// CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR5]]) +global int *global FileVar5; +// CHECK-DAG: ![[FILEVAR6:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar6", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true) +// CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR6]]) +constant int *global FileVar6; +// CHECK-DAG: ![[FILEVAR7:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar7", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true) +// CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR7]]) +local int *global FileVar7; +// CHECK-DAG: ![[FILEVAR8:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar8", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true) +// CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR8]]) +private int *global FileVar8; +// CHECK-DAG: ![[FILEVAR9:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar9", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true) +// CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR9]]) +int *global FileVar9; + +// CHECK-DAG: ![[FILEVAR10:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar10", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true) +// CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR10]]) +global int *constant FileVar10 = 0; +// CHECK-DAG: ![[FILEVAR11:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar11", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true) +// CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR11]]) +constant int *constant FileVar11 = 0; +// CHECK-DAG: ![[FILEVAR12:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar12", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true) +// CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR12]]) +local int *constant FileVar12 = 0; +// CHECK-DAG: ![[FILEVAR13:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar13", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true) +// CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR13]]) +private int *constant FileVar13 = 0; +// CHECK-DAG: ![[FILEVAR14:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar14", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true) +// CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR14]]) +int *constant FileVar14 = 0; + +kernel void kernel1( + // CHECK-DAG: ![[KERNELARG0:[0-9]+]] = !DILocalVariable(name: "KernelArg0", arg: {{[0-9]+}}, scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}) + // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(1)** {{.*}}, metadata ![[KERNELARG0]], metadata ![[PRIVATE]]), !dbg !{{[0-9]+}} + global int *KernelArg0, + // CHECK-DAG: ![[KERNELARG1:[0-9]+]] = !DILocalVariable(name: "KernelArg1", arg: {{[0-9]+}}, scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}) + // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(2)** {{.*}}, metadata ![[KERNELARG1]], metadata ![[PRIVATE]]), !dbg !{{[0-9]+}} + constant int *KernelArg1, + // CHECK-DAG: ![[KERNELARG2:[0-9]+]] = !DILocalVariable(name: "KernelArg2", arg: {{[0-9]+}}, scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}) + // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(3)** {{.*}}, metadata ![[KERNELARG2]], metadata ![[PRIVATE]]), !dbg !{{[0-9]+}} + local int *KernelArg2) { + private int *Tmp0; + int *Tmp1; + + // CHECK-DAG: ![[FUNCVAR0:[0-9]+]] = !DILocalVariable(name: "FuncVar0", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}) + // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(1)** {{.*}}, metadata ![[FUNCVAR0]], metadata ![[PRIVATE]]), !dbg !{{[0-9]+}} + global int *FuncVar0 = KernelArg0; + // CHECK-DAG: ![[FUNCVAR1:[0-9]+]] = !DILocalVariable(name: "FuncVar1", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}) + // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(2)** {{.*}}, metadata ![[FUNCVAR1]], metadata ![[PRIVATE]]), !dbg !{{[0-9]+}} + constant int *FuncVar1 = KernelArg1; + // CHECK-DAG: ![[FUNCVAR2:[0-9]+]] = !DILocalVariable(name: "FuncVar2", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}) + // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(3)** {{.*}}, metadata ![[FUNCVAR2]], metadata ![[PRIVATE]]), !dbg !{{[0-9]+}} + local int *FuncVar2 = KernelArg2; + // CHECK-DAG: ![[FUNCVAR3:[0-9]+]] = !DILocalVariable(name: "FuncVar3", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}) + // CHECK-DAG: call void @llvm.dbg.declare(metadata i32** {{.*}}, metadata ![[FUNCVAR3]], metadata ![[PRIVATE]]), !dbg !{{[0-9]+}} + private int *FuncVar3 = Tmp0; + // CHECK-DAG: ![[FUNCVAR4:[0-9]+]] = !DILocalVariable(name: "FuncVar4", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}) + // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(4)** {{.*}}, metadata ![[FUNCVAR4]], metadata ![[PRIVATE]]), !dbg !{{[0-9]+}} + int *FuncVar4 = Tmp1; + + // CHECK-DAG: ![[FUNCVAR5:[0-9]+]] = !DILocalVariable(name: "FuncVar5", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}) + // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(1)** {{.*}}, metadata ![[FUNCVAR5]], metadata ![[NONE]]), !dbg !{{[0-9]+}} + global int *constant FuncVar5 = KernelArg0; + // CHECK-DAG: ![[FUNCVAR6:[0-9]+]] = !DILocalVariable(name: "FuncVar6", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}) + // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(2)** {{.*}}, metadata ![[FUNCVAR6]], metadata ![[NONE]]), !dbg !{{[0-9]+}} + constant int *constant FuncVar6 = KernelArg1; + // CHECK-DAG: ![[FUNCVAR7:[0-9]+]] = !DILocalVariable(name: "FuncVar7", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}) + // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(3)** {{.*}}, metadata ![[FUNCVAR7]], metadata ![[NONE]]), !dbg !{{[0-9]+}} + local int *constant FuncVar7 = KernelArg2; + // CHECK-DAG: ![[FUNCVAR8:[0-9]+]] = !DILocalVariable(name: "FuncVar8", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}) + // CHECK-DAG: call void @llvm.dbg.declare(metadata i32** {{.*}}, metadata ![[FUNCVAR8]], metadata ![[NONE]]), !dbg !{{[0-9]+}} + private int *constant FuncVar8 = Tmp0; + // CHECK-DAG: ![[FUNCVAR9:[0-9]+]] = !DILocalVariable(name: "FuncVar9", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}) + // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(4)** {{.*}}, metadata ![[FUNCVAR9]], metadata ![[NONE]]), !dbg !{{[0-9]+}} + int *constant FuncVar9 = Tmp1; + + // CHECK-DAG: ![[FUNCVAR10:[0-9]+]] = distinct !DIGlobalVariable(name: "FuncVar10", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: true, isDefinition: true) + // CHECK-DAG: !DIGlobalVariableExpression(var: ![[FUNCVAR10]], expr: ![[LOCAL]]) + global int *local FuncVar10; FuncVar10 = KernelArg0; + // CHECK-DAG: ![[FUNCVAR11:[0-9]+]] = distinct !DIGlobalVariable(name: "FuncVar11", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: true, isDefinition: true) + // CHECK-DAG: !DIGlobalVariableExpression(var: ![[FUNCVAR11]], expr: ![[LOCAL]]) + constant int *local FuncVar11; FuncVar11 = KernelArg1; + // CHECK-DAG: ![[FUNCVAR12:[0-9]+]] = distinct !DIGlobalVariable(name: "FuncVar12", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: true, isDefinition: true) + // CHECK-DAG: !DIGlobalVariableExpression(var: ![[FUNCVAR12]], expr: ![[LOCAL]]) + local int *local FuncVar12; FuncVar12 = KernelArg2; + // CHECK-DAG: ![[FUNCVAR13:[0-9]+]] = distinct !DIGlobalVariable(name: "FuncVar13", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: true, isDefinition: true) + // CHECK-DAG: !DIGlobalVariableExpression(var: ![[FUNCVAR13]], expr: ![[LOCAL]]) + private int *local FuncVar13; FuncVar13 = Tmp0; + // CHECK-DAG: ![[FUNCVAR14:[0-9]+]] = distinct !DIGlobalVariable(name: "FuncVar14", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: true, isDefinition: true) + // CHECK-DAG: !DIGlobalVariableExpression(var: ![[FUNCVAR14]], expr: ![[LOCAL]]) + int *local FuncVar14; FuncVar14 = Tmp1; + + // CHECK-DAG: ![[FUNCVAR15:[0-9]+]] = !DILocalVariable(name: "FuncVar15", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}) + // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(1)** {{.*}}, metadata ![[FUNCVAR15]], metadata ![[PRIVATE]]), !dbg !{{[0-9]+}} + global int *private FuncVar15 = KernelArg0; + // CHECK-DAG: ![[FUNCVAR16:[0-9]+]] = !DILocalVariable(name: "FuncVar16", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}) + // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(2)** {{.*}}, metadata ![[FUNCVAR16]], metadata ![[PRIVATE]]), !dbg !{{[0-9]+}} + constant int *private FuncVar16 = KernelArg1; + // CHECK-DAG: ![[FUNCVAR17:[0-9]+]] = !DILocalVariable(name: "FuncVar17", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}) + // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(3)** {{.*}}, metadata ![[FUNCVAR17]], metadata ![[PRIVATE]]), !dbg !{{[0-9]+}} + local int *private FuncVar17 = KernelArg2; + // CHECK-DAG: ![[FUNCVAR18:[0-9]+]] = !DILocalVariable(name: "FuncVar18", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}) + // CHECK-DAG: call void @llvm.dbg.declare(metadata i32** {{.*}}, metadata ![[FUNCVAR18]], metadata ![[PRIVATE]]), !dbg !{{[0-9]+}} + private int *private FuncVar18 = Tmp0; + // CHECK-DAG: ![[FUNCVAR19:[0-9]+]] = !DILocalVariable(name: "FuncVar19", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}) + // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(4)** {{.*}}, metadata ![[FUNCVAR19]], metadata ![[PRIVATE]]), !dbg !{{[0-9]+}} + int *private FuncVar19 = Tmp1; +} diff --git a/test/CodeGenOpenCL/amdgpu-env-amdgiz.cl b/test/CodeGenOpenCL/amdgpu-env-amdgiz.cl new file mode 100644 index 000000000000..4cb8d95fa098 --- /dev/null +++ b/test/CodeGenOpenCL/amdgpu-env-amdgiz.cl @@ -0,0 +1,9 @@ +// RUN: %clang_cc1 %s -O0 -triple amdgcn -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -O0 -triple amdgcn---opencl -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -O0 -triple amdgcn---amdgiz -emit-llvm -o - | FileCheck -check-prefix=GIZ %s +// RUN: %clang_cc1 %s -O0 -triple amdgcn---amdgizcl -emit-llvm -o - | FileCheck -check-prefix=GIZ %s + +// CHECK: target datalayout = "e-p:32:32-p1:64:64-p2:64:64-p3:32:32-p4:64:64-p5:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64" +// GIZ: target datalayout = "e-p:64:64-p1:64:64-p2:64:64-p3:32:32-p4:32:32-p5:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-A5" +void foo(void) {} + diff --git a/test/CodeGenOpenCL/amdgpu-nullptr.cl b/test/CodeGenOpenCL/amdgpu-nullptr.cl index 2d36fa90a6ab..37c7469aa180 100644 --- a/test/CodeGenOpenCL/amdgpu-nullptr.cl +++ b/test/CodeGenOpenCL/amdgpu-nullptr.cl @@ -20,7 +20,7 @@ typedef struct { // Test 0 as initializer. -// CHECK: @private_p = local_unnamed_addr addrspace(1) global i8* addrspacecast (i8 addrspace(4)* null to i8*), align 4 +// CHECK: @private_p = local_unnamed_addr addrspace(1) global i8* null, 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 @@ -37,7 +37,7 @@ 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 +// CHECK: @private_p_NULL = local_unnamed_addr addrspace(1) global i8* null, 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 @@ -58,38 +58,55 @@ generic char *generic_p_NULL = NULL; // 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 +// CHECK: @fold_priv = local_unnamed_addr addrspace(1) global i16* null, 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 +// CHECK: @fold_priv_arith = local_unnamed_addr addrspace(1) global i8* inttoptr (i32 10 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 +// CHECK: @fold_int = local_unnamed_addr addrspace(1) global i32 14, 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 +// CHECK: @fold_int2 = local_unnamed_addr addrspace(1) global i32 13, align 4 int fold_int2 = (int) ((private void*)0 + 13); -// CHECK: @fold_int3 = local_unnamed_addr addrspace(1) global i32 -1, align 4 +// CHECK: @fold_int3 = local_unnamed_addr addrspace(1) global i32 0, align 4 int fold_int3 = (int) ((private int*)0); -// CHECK: @fold_int4 = local_unnamed_addr addrspace(1) global i32 7, align 4 +// CHECK: @fold_int4 = local_unnamed_addr addrspace(1) global i32 8, align 4 int fold_int4 = (int) &((private int*)0)[2]; -// CHECK: @fold_int5 = local_unnamed_addr addrspace(1) global i32 3, align 4 +// CHECK: @fold_int5 = local_unnamed_addr addrspace(1) global i32 4, align 4 int fold_int5 = (int) &((private StructTy1*)0)->p2; + +// CHECK: @fold_int_local = local_unnamed_addr addrspace(1) global i32 13, align 4 +int fold_int_local = (int)(local void*)(generic char*)(global int*)0 + 14; + +// CHECK: @fold_int2_local = local_unnamed_addr addrspace(1) global i32 12, align 4 +int fold_int2_local = (int) ((local void*)0 + 13); + +// CHECK: @fold_int3_local = local_unnamed_addr addrspace(1) global i32 -1, align 4 +int fold_int3_local = (int) ((local int*)0); + +// CHECK: @fold_int4_local = local_unnamed_addr addrspace(1) global i32 7, align 4 +int fold_int4_local = (int) &((local int*)0)[2]; + +// CHECK: @fold_int5_local = local_unnamed_addr addrspace(1) global i32 3, align 4 +int fold_int5_local = (int) &((local 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 +// NOOPT: @test_static_var_private.sp1 = internal addrspace(1) global i8* null, align 4 +// NOOPT: @test_static_var_private.sp2 = internal addrspace(1) global i8* null, align 4 +// NOOPT: @test_static_var_private.sp3 = internal addrspace(1) global i8* null, align 4 +// NOOPT: @test_static_var_private.sp4 = internal addrspace(1) global i8* null, align 4 +// NOOPT: @test_static_var_private.sp5 = internal addrspace(1) global i8* null, align 4 +// NOOPT: @test_static_var_private.SS1 = internal addrspace(1) global %struct.StructTy1 { i8* null, 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_private.SS2 = internal addrspace(1) global %struct.StructTy2 zeroinitializer, align 4 -void test_static_var(void) { +void test_static_var_private(void) { static private char *sp1 = 0; static private char *sp2 = NULL; static private char *sp3; @@ -100,18 +117,35 @@ void test_static_var(void) { static StructTy2 SS2; } +// NOOPT: @test_static_var_local.sp1 = internal addrspace(1) global i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), align 4 +// NOOPT: @test_static_var_local.sp2 = internal addrspace(1) global i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), align 4 +// NOOPT: @test_static_var_local.sp3 = internal addrspace(1) global i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), align 4 +// NOOPT: @test_static_var_local.sp4 = internal addrspace(1) global i8 addrspace(3)* null, align 4 +// NOOPT: @test_static_var_local.sp5 = internal addrspace(1) global i8 addrspace(3)* null, align 4 +// NOOPT: @test_static_var_local.SS1 = internal addrspace(1) global %struct.StructTy1 { i8* null, 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_local.SS2 = internal addrspace(1) global %struct.StructTy2 zeroinitializer, align 4 +void test_static_var_local(void) { + static local char *sp1 = 0; + static local char *sp2 = NULL; + static local char *sp3; + static local char *sp4 = (local char*)((void)0, 0); + const int x = 0; + static local char *sp5 = (local 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-LABEL: @test_func_scope_var_private( +// NOOPT: store i8* null, i8** %sp1, align 4 +// NOOPT: store i8* null, 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: call void @llvm.memcpy.p0i8.p2i8.i64(i8* %[[SS1]], i8 addrspace(2)* bitcast (%struct.StructTy1 addrspace(2)* @test_func_scope_var_private.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) { +void test_func_scope_var_private(void) { private char *sp1 = 0; private char *sp2 = NULL; private char *sp3 = (private char*)((void)0, 0); @@ -121,13 +155,34 @@ void test_func_scope_var(void) { StructTy2 SS2 = {0, 0, 0}; } +// Test function-scope variable initialization. +// NOOPT-LABEL: @test_func_scope_var_local( +// NOOPT: store i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), i8 addrspace(3)** %sp1, align 4 +// NOOPT: store i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), i8 addrspace(3)** %sp2, align 4 +// NOOPT: store i8 addrspace(3)* null, i8 addrspace(3)** %sp3, align 4 +// NOOPT: store i8 addrspace(3)* null, i8 addrspace(3)** %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_local.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_local(void) { + local char *sp1 = 0; + local char *sp2 = NULL; + local char *sp3 = (local char*)((void)0, 0); + const int x = 0; + local char *sp4 = (local 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 +// CHECK: @p1 = common local_unnamed_addr addrspace(1) global i8* null, 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 @@ -144,14 +199,14 @@ 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 +// CHECK: @S1 = weak local_unnamed_addr addrspace(1) global %struct.StructTy1 { i8* null, 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 +// CHECK: @A1 = weak local_unnamed_addr addrspace(1) global [2 x %struct.StructTy1] [%struct.StructTy1 { i8* null, 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* null, 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 @@ -160,7 +215,7 @@ StructTy2 A2[2]; // Test comparison with 0. // CHECK-LABEL: cmp_private -// CHECK: icmp eq i8* %p, addrspacecast (i8 addrspace(4)* null to i8*) +// CHECK: icmp eq i8* %p, null void cmp_private(private char* p) { if (p != 0) *p = 0; @@ -199,7 +254,7 @@ void cmp_generic(generic char* p) { // Test comparison with NULL. // CHECK-LABEL: cmp_NULL_private -// CHECK: icmp eq i8* %p, addrspacecast (i8 addrspace(4)* null to i8*) +// CHECK: icmp eq i8* %p, null void cmp_NULL_private(private char* p) { if (p != NULL) *p = 0; @@ -237,7 +292,7 @@ void cmp_NULL_generic(generic char* p) { // 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* null, 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 @@ -256,7 +311,7 @@ void test_storage_null_pointer(private char** arg_private, // 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* null, 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 @@ -281,8 +336,8 @@ void test_pass_null_pointer_arg_calee(private char* arg_private, 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) +// CHECK: call void @test_pass_null_pointer_arg_calee(i8* null, 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* null, 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); @@ -296,8 +351,8 @@ void test_cast_null_pointer_to_sizet_calee(size_t arg_private, 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) +// CHECK: call void @test_cast_null_pointer_to_sizet_calee(i64 0, 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 0, 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), @@ -409,7 +464,7 @@ 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*) +// CHECK: icmp eq i8* %p, null void cast_bool_private(private char* p) { if (p) *p = 0; @@ -453,10 +508,9 @@ typedef struct { 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) { +// CHECK-LABEL: test_memset_private +// CHECK: call void @llvm.memset.p0i8.i64(i8* nonnull {{.*}}, i8 0, i64 40, i32 8, i1 false) +StructTy3 test_memset_private(void) { StructTy3 S3 = {0, 0, 0, 0, 0}; return S3; } @@ -464,22 +518,34 @@ StructTy3 test_memset(void) { // 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) { +// CHECK-LABEL: test_cast_0_to_local_ptr +// CHECK: ret i32 addrspace(3)* addrspacecast (i32 addrspace(4)* null to i32 addrspace(3)*) +local int* test_cast_0_to_local_ptr(void) { + return (local int*)0; +} + +// CHECK-LABEL: test_cast_0_to_private_ptr +// CHECK: ret i32* null +private int* test_cast_0_to_private_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. +// zero value. -// CHECK-LABEL: test_cast_int_to_ptr1 +// CHECK-LABEL: test_cast_int_to_ptr1_private // CHECK: ret i32* null -private int* test_cast_int_to_ptr1(void) { +private int* test_cast_int_to_ptr1_private(void) { return (private int*)((void)0, 0); } +// CHECK-LABEL: test_cast_int_to_ptr1_local + // CHECK: ret i32 addrspace(3)* null +local int* test_cast_int_to_ptr1_local(void) { + return (local int*)((void)0, 0); +} + // CHECK-LABEL: test_cast_int_to_ptr2 // CHECK: ret i32* null private int* test_cast_int_to_ptr2(void) { @@ -500,15 +566,25 @@ 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-LABEL: test_not_private_ptr +// CHECK: %[[lnot:.*]] = icmp eq i8* %p, null +// CHECK: %[[lnot_ext:.*]] = zext i1 %[[lnot]] to i32 +// CHECK: ret i32 %[[lnot_ext]] +int test_not_private_ptr(private char* p) { + return !p; +} + +// CHECK-LABEL: test_not_local_ptr +// CHECK: %[[lnot:.*]] = icmp eq i8 addrspace(3)* %p, addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*) // CHECK: %[[lnot_ext:.*]] = zext i1 %[[lnot]] to i32 // CHECK: ret i32 %[[lnot_ext]] -int test_not_ptr(private char* p) { +int test_not_local_ptr(local char* p) { return !p; } + + // CHECK-LABEL: test_and_ptr -// CHECK: %[[tobool:.*]] = icmp ne i8* %p1, addrspacecast (i8 addrspace(4)* null to i8*) +// CHECK: %[[tobool:.*]] = icmp ne i8* %p1, null // 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 @@ -518,17 +594,31 @@ int test_and_ptr(private char* p1, local char* p2) { } // Test folding of null pointer in function scope. -// NOOPT-LABEL: test_fold +// NOOPT-LABEL: test_fold_private // 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: %{{.*}} = add nsw i64 %{{.*}}, 0 // NOOPT: %{{.*}} = sub nsw i64 %{{.*}}, 1 void test_fold_callee(void); -void test_fold(void) { +void test_fold_private(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); } + +// NOOPT-LABEL: test_fold_local +// 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 addrspace(3)* addrspacecast (i32 addrspace(4)* null to i32 addrspace(3)*) to i32) to i64) +// NOOPT: %{{.*}} = sub nsw i64 %{{.*}}, 1 +void test_fold_local(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(), (local int*)(generic char*)(global short*)0); + x = x - (int)((local int*)0 == (local int*)(generic char*)0); +} diff --git a/test/CodeGenOpenCL/blocks.cl b/test/CodeGenOpenCL/blocks.cl new file mode 100644 index 000000000000..5f0cceaf6cdb --- /dev/null +++ b/test/CodeGenOpenCL/blocks.cl @@ -0,0 +1,17 @@ +// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -O0 -triple spir-unknown-unknown | FileCheck -check-prefix=GENERIC -check-prefix=COMMON %s +// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -O0 -triple amdgcn-amd-amdhsa-opencl | FileCheck -check-prefix=AMD -check-prefix=COMMON %s + +// Checking for null instead of @__NSConcreteGlobalBlock symbol +// COMMON: @__block_literal_global = internal addrspace(1) constant { i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* } { i8** null +void (^block_A)(local void *) = ^(local void *a) { + return; +}; + +void foo(){ + int i; +// Checking for null instead of @_NSConcreteStackBlock symbol +// COMMON: store i8* null, i8** %block.isa + int (^ block_B)(void) = ^{ + return i; + }; +} diff --git a/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl b/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl new file mode 100644 index 000000000000..333b610f81bd --- /dev/null +++ b/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl @@ -0,0 +1,11 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -S -emit-llvm -o - %s | FileCheck %s + +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +// CHECK-LABEL: @test_fmed3_f16 +// CHECK: call half @llvm.amdgcn.fmed3.f16(half %a, half %b, half %c) +void test_fmed3_f16(global half* out, half a, half b, half c) +{ + *out = __builtin_amdgcn_fmed3h(a, b, c); +} diff --git a/test/CodeGenOpenCL/builtins-amdgcn-vi.cl b/test/CodeGenOpenCL/builtins-amdgcn-vi.cl index 0c88fdfb724d..1dad67491832 100644 --- a/test/CodeGenOpenCL/builtins-amdgcn-vi.cl +++ b/test/CodeGenOpenCL/builtins-amdgcn-vi.cl @@ -81,3 +81,11 @@ void test_s_memrealtime(global ulong* out) { *out = __builtin_amdgcn_s_memrealtime(); } + +// CHECK-LABEL: @test_mov_dpp +// CHECK: call i32 @llvm.amdgcn.mov.dpp.i32(i32 %src, i32 0, i32 0, i32 0, i1 false) +void test_mov_dpp(global int* out, int src) +{ + *out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0, false); +} + diff --git a/test/CodeGenOpenCL/builtins-amdgcn.cl b/test/CodeGenOpenCL/builtins-amdgcn.cl index ce566b6e876a..a19ce2f1ed59 100644 --- a/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -235,6 +235,34 @@ void test_ds_swizzle(global int* out, int a) *out = __builtin_amdgcn_ds_swizzle(a, 32); } +// CHECK-LABEL: @test_ds_permute +// CHECK: call i32 @llvm.amdgcn.ds.permute(i32 %a, i32 %b) +void test_ds_permute(global int* out, int a, int b) +{ + out[0] = __builtin_amdgcn_ds_permute(a, b); +} + +// CHECK-LABEL: @test_ds_bpermute +// CHECK: call i32 @llvm.amdgcn.ds.bpermute(i32 %a, i32 %b) +void test_ds_bpermute(global int* out, int a, int b) +{ + *out = __builtin_amdgcn_ds_bpermute(a, b); +} + +// CHECK-LABEL: @test_readfirstlane +// CHECK: call i32 @llvm.amdgcn.readfirstlane(i32 %a) +void test_readfirstlane(global int* out, int a) +{ + *out = __builtin_amdgcn_readfirstlane(a); +} + +// CHECK-LABEL: @test_readlane +// CHECK: call i32 @llvm.amdgcn.readlane(i32 %a, i32 %b) +void test_readlane(global int* out, int a, int b) +{ + *out = __builtin_amdgcn_readlane(a, b); +} + // 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) @@ -263,6 +291,55 @@ void test_class_f64(global double* out, double a, int b) *out = __builtin_amdgcn_class(a, b); } +// CHECK-LABEL: @test_buffer_wbinvl1 +// CHECK: call void @llvm.amdgcn.buffer.wbinvl1( +void test_buffer_wbinvl1() +{ + __builtin_amdgcn_buffer_wbinvl1(); +} + +// CHECK-LABEL: @test_s_dcache_inv +// CHECK: call void @llvm.amdgcn.s.dcache.inv( +void test_s_dcache_inv() +{ + __builtin_amdgcn_s_dcache_inv(); +} + +// CHECK-LABEL: @test_s_waitcnt +// CHECK: call void @llvm.amdgcn.s.waitcnt( +void test_s_waitcnt() +{ + __builtin_amdgcn_s_waitcnt(0); +} + +// CHECK-LABEL: @test_s_sendmsg +// CHECK: call void @llvm.amdgcn.s.sendmsg( +void test_s_sendmsg() +{ + __builtin_amdgcn_s_sendmsg(1, 0); +} + +// CHECK-LABEL: @test_s_sendmsg_var +// CHECK: call void @llvm.amdgcn.s.sendmsg( +void test_s_sendmsg_var(int in) +{ + __builtin_amdgcn_s_sendmsg(1, in); +} + +// CHECK-LABEL: @test_s_sendmsghalt +// CHECK: call void @llvm.amdgcn.s.sendmsghalt( +void test_s_sendmsghalt() +{ + __builtin_amdgcn_s_sendmsghalt(1, 0); +} + +// CHECK-LABEL: @test_s_sendmsghalt +// CHECK: call void @llvm.amdgcn.s.sendmsghalt( +void test_s_sendmsghalt_var(int in) +{ + __builtin_amdgcn_s_sendmsghalt(1, in); +} + // CHECK-LABEL: @test_s_barrier // CHECK: call void @llvm.amdgcn.s.barrier( void test_s_barrier() @@ -371,6 +448,17 @@ void test_get_group_id(int d, global int *out) } } +// CHECK-LABEL: @test_s_getreg( +// CHECK: tail call i32 @llvm.amdgcn.s.getreg(i32 0) +// CHECK: tail call i32 @llvm.amdgcn.s.getreg(i32 1) +// CHECK: tail call i32 @llvm.amdgcn.s.getreg(i32 65535) +void test_s_getreg(volatile global uint *out) +{ + *out = __builtin_amdgcn_s_getreg(0); + *out = __builtin_amdgcn_s_getreg(1); + *out = __builtin_amdgcn_s_getreg(65535); +} + // CHECK-LABEL: @test_get_local_id( // CHECK: tail call i32 @llvm.amdgcn.workitem.id.x(), !range [[WI_RANGE:![0-9]*]] // CHECK: tail call i32 @llvm.amdgcn.workitem.id.y(), !range [[WI_RANGE]] @@ -385,6 +473,13 @@ void test_get_local_id(int d, global int *out) } } +// CHECK-LABEL: @test_fmed3_f32 +// CHECK: call float @llvm.amdgcn.fmed3.f32( +void test_fmed3_f32(global float* out, float a, float b, float c) +{ + *out = __builtin_amdgcn_fmed3f(a, b, c); +} + // CHECK-DAG: [[WI_RANGE]] = !{i32 0, i32 1024} // CHECK-DAG: attributes #[[NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly } // CHECK-DAG: attributes #[[READ_EXEC_ATTRS]] = { convergent } diff --git a/test/CodeGenOpenCL/cl20-device-side-enqueue.cl b/test/CodeGenOpenCL/cl20-device-side-enqueue.cl index add13b9e3fc4..def290661534 100644 --- a/test/CodeGenOpenCL/cl20-device-side-enqueue.cl +++ b/test/CodeGenOpenCL/cl20-device-side-enqueue.cl @@ -2,9 +2,10 @@ // 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 *); +typedef struct {int a;} ndrange_t; // N.B. The check here only exists to set BL_GLOBAL -// COMMON: @block_G = {{.*}}bitcast ([[BL_GLOBAL:[^@]+@__block_literal_global(\.[0-9]+)?]] +// COMMON: @block_G = addrspace(1) constant void (i8 addrspace(3)*) addrspace(4)* addrspacecast (void (i8 addrspace(3)*) addrspace(1)* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* } addrspace(1)* [[BL_GLOBAL:@__block_literal_global(\.[0-9]+)?]] to void (i8 addrspace(3)*) addrspace(1)*) to void (i8 addrspace(3)*) addrspace(4)*) const bl_t block_G = (bl_t) ^ (local void *a) {}; kernel void device_side_enqueue(global int *a, global int *b, int i) { @@ -12,7 +13,7 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) { queue_t default_queue; // COMMON: %flags = alloca i32 unsigned flags = 0; - // COMMON: %ndrange = alloca %opencl.ndrange_t* + // COMMON: %ndrange = alloca %struct.ndrange_t ndrange_t ndrange; // COMMON: %clk_event = alloca %opencl.clk_event_t* clk_event_t clk_event; @@ -21,12 +22,12 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) { // COMMON: %event_wait_list2 = alloca [1 x %opencl.clk_event_t*] clk_event_t event_wait_list2[] = {clk_event}; + // COMMON: [[NDR:%[a-z0-9]+]] = alloca %struct.ndrange_t, align 4 // 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]]) + // COMMON: [[BL_I8:%[0-9]+]] = addrspacecast void ()* [[BL]] to i8 addrspace(4)* + // COMMON: call i32 @__enqueue_kernel_basic(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* byval [[NDR]]{{(.[0-9]+)?}}, i8 addrspace(4)* [[BL_I8]]) enqueue_kernel(default_queue, flags, ndrange, ^(void) { a[i] = b[i]; @@ -34,12 +35,11 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) { // 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]]) + // COMMON: [[BL_I8:%[0-9]+]] = addrspacecast void ()* [[BL]] to i8 addrspace(4)* + // COMMON: call i32 @__enqueue_kernel_basic_events(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* {{.*}}, i32 2, %opencl.clk_event_t{{.*}}* addrspace(4)* [[WAIT_EVNT]], %opencl.clk_event_t{{.*}}* addrspace(4)* [[EVNT]], i8 addrspace(4)* [[BL_I8]]) enqueue_kernel(default_queue, flags, ndrange, 2, &event_wait_list, &clk_event, ^(void) { a[i] = b[i]; @@ -47,9 +47,8 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) { // 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) + // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %struct.ndrange_t*, i8 addrspace(4)*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* [[NDR]]{{(.[0-9]+)?}}, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* } addrspace(1)* @__block_literal_global{{(.[0-9]+)?}} to i8 addrspace(1)*) to i8 addrspace(4)*), i32 1, i32 256) + // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %struct.ndrange_t*, i8 addrspace(4)*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* [[NDR]]{{(.[0-9]+)?}}, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* } addrspace(1)* @__block_literal_global{{(.[0-9]+)?}} to i8 addrspace(1)*) to i8 addrspace(4)*), i32 1, i64 256) enqueue_kernel(default_queue, flags, ndrange, ^(local void *p) { return; @@ -58,11 +57,10 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) { char c; // 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]]) + // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %struct.ndrange_t*, i8 addrspace(4)*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* [[NDR]]{{(.[0-9]+)?}}, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* } addrspace(1)* @__block_literal_global{{(.[0-9]+)?}} to i8 addrspace(1)*) to i8 addrspace(4)*), 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]]) + // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %struct.ndrange_t*, i8 addrspace(4)*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* [[NDR]]{{(.[0-9]+)?}}, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* } addrspace(1)* @__block_literal_global{{(.[0-9]+)?}} to i8 addrspace(1)*) to i8 addrspace(4)*), i32 1, i64 [[SIZE]]) enqueue_kernel(default_queue, flags, ndrange, ^(local void *p) { return; @@ -71,12 +69,11 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) { // 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) + // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %struct.ndrange_t*, i32, %opencl.clk_event_t{{.*}}* addrspace(4)*, %opencl.clk_event_t{{.*}}* addrspace(4)*, i8 addrspace(4)*, i32, ...) @__enqueue_kernel_events_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* {{.*}}, i32 2, %opencl.clk_event_t{{.*}} [[WAIT_EVNT]], %opencl.clk_event_t{{.*}} [[EVNT]], i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* } addrspace(1)* @__block_literal_global{{(.[0-9]+)?}} to i8 addrspace(1)*) to i8 addrspace(4)*), i32 1, i32 256) + // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %struct.ndrange_t*, i32, %opencl.clk_event_t{{.*}}* addrspace(4)*, %opencl.clk_event_t{{.*}}* addrspace(4)*, i8 addrspace(4)*, i32, ...) @__enqueue_kernel_events_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* {{.*}}, i32 2, %opencl.clk_event_t{{.*}} [[WAIT_EVNT]], %opencl.clk_event_t{{.*}} [[EVNT]], i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* } addrspace(1)* @__block_literal_global{{(.[0-9]+)?}} to i8 addrspace(1)*) to i8 addrspace(4)*), i32 1, i64 256) enqueue_kernel(default_queue, flags, ndrange, 2, event_wait_list2, &clk_event, ^(local void *p) { return; @@ -85,14 +82,13 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) { // 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]]) + // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %struct.ndrange_t*, i32, %opencl.clk_event_t{{.*}}* addrspace(4)*, %opencl.clk_event_t{{.*}}* addrspace(4)*, i8 addrspace(4)*, i32, ...) @__enqueue_kernel_events_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* {{.*}}, i32 2, %opencl.clk_event_t{{.*}}* addrspace(4)* [[WAIT_EVNT]], %opencl.clk_event_t{{.*}}* addrspace(4)* [[EVNT]], i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* } addrspace(1)* @__block_literal_global{{(.[0-9]+)?}} to i8 addrspace(1)*) to i8 addrspace(4)*), 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]]) + // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %struct.ndrange_t*, i32, %opencl.clk_event_t{{.*}}* addrspace(4)*, %opencl.clk_event_t{{.*}}* addrspace(4)*, i8 addrspace(4)*, i32, ...) @__enqueue_kernel_events_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* {{.*}}, i32 2, %opencl.clk_event_t{{.*}}* addrspace(4)* [[WAIT_EVNT]], %opencl.clk_event_t{{.*}}* addrspace(4)* [[EVNT]], i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* } addrspace(1)* @__block_literal_global{{(.[0-9]+)?}} to i8 addrspace(1)*) to i8 addrspace(4)*), i32 1, i64 [[SIZE]]) enqueue_kernel(default_queue, flags, ndrange, 2, event_wait_list2, &clk_event, ^(local void *p) { return; @@ -102,11 +98,10 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) { 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]]) + // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %struct.ndrange_t*, i8 addrspace(4)*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* [[NDR]]{{(.[0-9]+)?}}, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* } addrspace(1)* @__block_literal_global{{(.[0-9]+)?}} to i8 addrspace(1)*) to i8 addrspace(4)*), 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]]) + // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %struct.ndrange_t*, i8 addrspace(4)*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* [[NDR]]{{(.[0-9]+)?}}, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* } addrspace(1)* @__block_literal_global{{(.[0-9]+)?}} to i8 addrspace(1)*) to i8 addrspace(4)*), i32 1, i64 [[SIZE]]) enqueue_kernel(default_queue, flags, ndrange, ^(local void *p) { return; @@ -115,9 +110,8 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) { // 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) + // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %struct.ndrange_t*, i8 addrspace(4)*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* [[NDR]]{{(.[0-9]+)?}}, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* } addrspace(1)* @__block_literal_global{{(.[0-9]+)?}} to i8 addrspace(1)*) to i8 addrspace(4)*), i32 1, i32 0) + // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %struct.ndrange_t*, i8 addrspace(4)*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* [[NDR]]{{(.[0-9]+)?}}, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* } addrspace(1)* @__block_literal_global{{(.[0-9]+)?}} to i8 addrspace(1)*) to i8 addrspace(4)*), i32 1, i64 4294967296) enqueue_kernel(default_queue, flags, ndrange, ^(local void *p) { return; @@ -126,22 +120,22 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) { // 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 + // COMMON: store void () addrspace(4)* addrspacecast (void () addrspace(1)* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* } addrspace(1)* [[BL_A:@__block_literal_global(\.[0-9]+)?]] to void () addrspace(1)*) to void () addrspace(4)*), void () addrspace(4)** %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 + // COMMON: store void (i8 addrspace(3)*) addrspace(4)* addrspacecast (void (i8 addrspace(3)*) addrspace(1)* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* } addrspace(1)* [[BL_B:@__block_literal_global(\.[0-9]+)?]] to void (i8 addrspace(3)*) addrspace(1)*) to void (i8 addrspace(3)*) addrspace(4)*), void (i8 addrspace(3)*) addrspace(4)** %block_B void (^const block_B)(local void *) = ^(local void *a) { return; }; - // COMMON: call i32 @__get_kernel_work_group_size_impl(i8* bitcast ([[BL_A]] to i8*)) + // COMMON: call i32 @__get_kernel_work_group_size_impl(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* } addrspace(1)* [[BL_A]] to i8 addrspace(1)*) to i8 addrspace(4)*)) unsigned size = get_kernel_work_group_size(block_A); - // COMMON: call i32 @__get_kernel_work_group_size_impl(i8* bitcast ([[BL_B]] to i8*)) + // COMMON: call i32 @__get_kernel_work_group_size_impl(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* } addrspace(1)* [[BL_B]] to i8 addrspace(1)*) to i8 addrspace(4)*)) size = get_kernel_work_group_size(block_B); - // COMMON: call i32 @__get_kernel_preferred_work_group_multiple_impl(i8* bitcast ([[BL_A]] to i8*)) + // COMMON: call i32 @__get_kernel_preferred_work_group_multiple_impl(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* } addrspace(1)* [[BL_A]] to i8 addrspace(1)*) to i8 addrspace(4)*)) size = get_kernel_preferred_work_group_size_multiple(block_A); - // COMMON: call i32 @__get_kernel_preferred_work_group_multiple_impl(i8* bitcast ([[BL_GLOBAL]] to i8*)) + // COMMON: call i32 @__get_kernel_preferred_work_group_multiple_impl(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* } addrspace(1)* [[BL_GLOBAL]] to i8 addrspace(1)*) to i8 addrspace(4)*)) size = get_kernel_preferred_work_group_size_multiple(block_G); } diff --git a/test/CodeGenOpenCL/denorms-are-zero.cl b/test/CodeGenOpenCL/denorms-are-zero.cl index f398218b4fd7..ca7969ad5f40 100644 --- a/test/CodeGenOpenCL/denorms-are-zero.cl +++ b/test/CodeGenOpenCL/denorms-are-zero.cl @@ -1,19 +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 +// RUN: %clang_cc1 -emit-llvm -target-feature +fp32-denormals -target-feature -fp64-fp16-denormals -cl-denorms-are-zero -o - -triple amdgcn--amdhsa -target-cpu fiji %s | FileCheck --check-prefix=CHECK-FEATURE %s // 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. +// explicitly set. amdgcn target always do not flush fp64 denormals. The control for fp64 and fp16 denormals is the same. // CHECK-DENORM-LABEL: define void @f() -// CHECK-DENORM: attributes #{{[0-9]*}} = {{{[^}]*}} "target-features"="{{[^"]*}}+fp64-denormals,{{[^"]*}}-fp32-denormals{{[^"]*}}" +// CHECK-DENORM: attributes #{{[0-9]*}} = {{{[^}]*}} "target-features"="{{[^"]*}}+fp64-fp16-denormals,{{[^"]*}}-fp32-denormals{{[^"]*}}" // CHECK-LABEL: define void @f() -// CHECK: attributes #{{[0-9]*}} = {{{[^}]*}} "target-features"="{{[^"]*}}+fp64-denormals,{{[^"]*}}-fp32-denormals{{[^"]*}}" +// CHECK: attributes #{{[0-9]*}} = {{{[^}]*}} "target-features"="{{[^"]*}}+fp64-fp16-denormals,{{[^"]*}}-fp32-denormals{{[^"]*}}" // CHECK-FEATURE-LABEL: define void @f() -// CHECK-FEATURE: attributes #{{[0-9]*}} = {{{[^}]*}} "target-features"="{{[^"]*}}+fp32-denormals,{{[^"]*}}-fp64-denormals{{[^"]*}}" +// CHECK-FEATURE: attributes #{{[0-9]*}} = {{{[^}]*}} "target-features"="{{[^"]*}}+fp32-denormals,{{[^"]*}}-fp64-fp16-denormals{{[^"]*}}" void f() {} diff --git a/test/CodeGenOpenCL/gfx9-fp32-denorms.cl b/test/CodeGenOpenCL/gfx9-fp32-denorms.cl new file mode 100644 index 000000000000..ccb4c6ddcf56 --- /dev/null +++ b/test/CodeGenOpenCL/gfx9-fp32-denorms.cl @@ -0,0 +1,13 @@ +// REQUIRES: amdgpu-registered-target + +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -S -emit-llvm -o - %s | FileCheck --check-prefix=DEFAULT %s +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -S -emit-llvm -o - -target-feature +fp32-denormals %s | FileCheck --check-prefix=FEATURE_FP32_DENORMALS_ON %s +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -S -emit-llvm -o - -target-feature -fp32-denormals %s | FileCheck --check-prefix=FEATURE_FP32_DENORMALS_OFF %s +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -S -emit-llvm -o - -cl-denorms-are-zero %s | FileCheck --check-prefix=OPT_DENORMS_ARE_ZERO %s + +// DEFAULT: +fp32-denormals +// FEATURE_FP32_DENORMALS_ON: +fp32-denormals +// FEATURE_FP32_DENORMALS_OFF: -fp32-denormals +// OPT_DENORMS_ARE_ZERO: -fp32-denormals + +kernel void gfx9_fp32_denorms() {} diff --git a/test/CodeGenOpenCL/kernel-arg-info.cl b/test/CodeGenOpenCL/kernel-arg-info.cl index 32017f85df8c..f1b2f45d7d97 100644 --- a/test/CodeGenOpenCL/kernel-arg-info.cl +++ b/test/CodeGenOpenCL/kernel-arg-info.cl @@ -2,7 +2,8 @@ // 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) { + volatile int anotherArg, __constant float * restrict Z, + __global volatile int * V, __global const int * C) { *X = Y + anotherArg; } // CHECK: define spir_kernel void @foo{{[^!]+}} @@ -60,11 +61,11 @@ kernel void foo5(myImage img1, write_only image1d_t img2) { // CHECK-NOT: !kernel_arg_name // ARGINFO: !kernel_arg_name ![[MD54:[0-9]+]] -// CHECK: ![[MD11]] = !{i32 1, i32 0, i32 0, i32 2} -// CHECK: ![[MD12]] = !{!"none", !"none", !"none", !"none"} -// CHECK: ![[MD13]] = !{!"int*", !"int", !"int", !"float*"} -// CHECK: ![[MD14]] = !{!"restrict", !"const", !"volatile", !"restrict const"} -// ARGINFO: ![[MD15]] = !{!"X", !"Y", !"anotherArg", !"Z"} +// CHECK: ![[MD11]] = !{i32 1, i32 0, i32 0, i32 2, i32 1, i32 1} +// CHECK: ![[MD12]] = !{!"none", !"none", !"none", !"none", !"none", !"none"} +// CHECK: ![[MD13]] = !{!"int*", !"int", !"int", !"float*", !"int*", !"int*"} +// CHECK: ![[MD14]] = !{!"restrict", !"", !"", !"restrict const", !"volatile", !"const"} +// ARGINFO: ![[MD15]] = !{!"X", !"Y", !"anotherArg", !"Z", !"V", !"C"} // 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"} diff --git a/test/CodeGenOpenCL/overload.cl b/test/CodeGenOpenCL/overload.cl new file mode 100644 index 000000000000..f182cb5fd2af --- /dev/null +++ b/test/CodeGenOpenCL/overload.cl @@ -0,0 +1,46 @@ +// RUN: %clang_cc1 -cl-std=CL2.0 -emit-llvm -o - -triple spir-unknown-unknown %s | FileCheck %s + +typedef short short4 __attribute__((ext_vector_type(4))); + +// CHECK-DAG: declare spir_func <4 x i16> @_Z5clampDv4_sS_S_(<4 x i16>, <4 x i16>, <4 x i16>) +short4 __attribute__ ((overloadable)) clamp(short4 x, short4 minval, short4 maxval); +// CHECK-DAG: declare spir_func <4 x i16> @_Z5clampDv4_sss(<4 x i16>, i16 signext, i16 signext) +short4 __attribute__ ((overloadable)) clamp(short4 x, short minval, short maxval); +void __attribute__((overloadable)) foo(global int *a, global int *b); +void __attribute__((overloadable)) foo(generic int *a, generic int *b); +void __attribute__((overloadable)) bar(generic int *global *a, generic int *global *b); +void __attribute__((overloadable)) bar(generic int *generic *a, generic int *generic *b); + +// Checking address space resolution +void kernel test1() { + global int *a; + global int *b; + generic int *c; + local int *d; + generic int *generic *gengen; + generic int *local *genloc; + generic int *global *genglob; + // CHECK-DAG: call spir_func void @_Z3fooPU3AS1iS0_(i32 addrspace(1)* undef, i32 addrspace(1)* undef) + foo(a, b); + // CHECK-DAG: call spir_func void @_Z3fooPU3AS4iS0_(i32 addrspace(4)* undef, i32 addrspace(4)* undef) + foo(b, c); + // CHECK-DAG: call spir_func void @_Z3fooPU3AS4iS0_(i32 addrspace(4)* undef, i32 addrspace(4)* undef) + foo(a, d); + + // CHECK-DAG: call spir_func void @_Z3barPU3AS4PU3AS4iS2_(i32 addrspace(4)* addrspace(4)* undef, i32 addrspace(4)* addrspace(4)* undef) + bar(gengen, genloc); + // CHECK-DAG: call spir_func void @_Z3barPU3AS4PU3AS4iS2_(i32 addrspace(4)* addrspace(4)* undef, i32 addrspace(4)* addrspace(4)* undef) + bar(gengen, genglob); + // CHECK-DAG: call spir_func void @_Z3barPU3AS1PU3AS4iS2_(i32 addrspace(4)* addrspace(1)* undef, i32 addrspace(4)* addrspace(1)* undef) + bar(genglob, genglob); +} + +// Checking vector vs scalar resolution +void kernel test2() { + short4 e0=0; + + // CHECK-DAG: call spir_func <4 x i16> @_Z5clampDv4_sss(<4 x i16> zeroinitializer, i16 signext 0, i16 signext 255) + clamp(e0, 0, 255); + // CHECK-DAG: call spir_func <4 x i16> @_Z5clampDv4_sS_S_(<4 x i16> zeroinitializer, <4 x i16> zeroinitializer, <4 x i16> zeroinitializer) + clamp(e0, e0, e0); +} diff --git a/test/CodeGenOpenCL/preserve_vec3.cl b/test/CodeGenOpenCL/preserve_vec3.cl new file mode 100644 index 000000000000..6efbbb3d5323 --- /dev/null +++ b/test/CodeGenOpenCL/preserve_vec3.cl @@ -0,0 +1,24 @@ +// RUN: %clang_cc1 %s -emit-llvm -o - -triple spir-unknown-unknown -fpreserve-vec3-type | FileCheck %s + +typedef float float3 __attribute__((ext_vector_type(3))); +typedef float float4 __attribute__((ext_vector_type(4))); + +void kernel foo(global float3 *a, global float3 *b) { + // CHECK: %[[LOAD_A:.*]] = load <3 x float>, <3 x float> addrspace(1)* %a + // CHECK: store <3 x float> %[[LOAD_A]], <3 x float> addrspace(1)* %b + *b = *a; +} + +void kernel float4_to_float3(global float3 *a, global float4 *b) { + // CHECK: %[[LOAD_A:.*]] = load <4 x float>, <4 x float> addrspace(1)* %b, align 16 + // CHECK: %[[ASTYPE:.*]] = shufflevector <4 x float> %[[LOAD_A]], <4 x float> undef, <3 x i32> <i32 0, i32 1, i32 2> + // CHECK: store <3 x float> %[[ASTYPE:.*]], <3 x float> addrspace(1)* %a, align 16 + *a = __builtin_astype(*b, float3); +} + +void kernel float3_to_float4(global float3 *a, global float4 *b) { + // CHECK: %[[LOAD_A:.*]] = load <3 x float>, <3 x float> addrspace(1)* %a, align 16 + // CHECK: %[[ASTYPE:.*]] = shufflevector <3 x float> %[[LOAD_A]], <3 x float> undef, <4 x i32> <i32 0, i32 1, i32 2, i32 undef> + // CHECK: store <4 x float> %[[ASTYPE:.*]], <4 x float> addrspace(1)* %b, align 16 + *b = __builtin_astype(*a, float4); +} diff --git a/test/CodeGenOpenCL/relaxed-fpmath.cl b/test/CodeGenOpenCL/relaxed-fpmath.cl index 6f302b1335b6..7676ee164ce4 100644 --- a/test/CodeGenOpenCL/relaxed-fpmath.cl +++ b/test/CodeGenOpenCL/relaxed-fpmath.cl @@ -2,33 +2,54 @@ // RUN: %clang_cc1 %s -emit-llvm -cl-fast-relaxed-math -o - | FileCheck %s -check-prefix=FAST // RUN: %clang_cc1 %s -emit-llvm -cl-finite-math-only -o - | FileCheck %s -check-prefix=FINITE // RUN: %clang_cc1 %s -emit-llvm -cl-unsafe-math-optimizations -o - | FileCheck %s -check-prefix=UNSAFE - -typedef __attribute__(( ext_vector_type(4) )) float float4; +// RUN: %clang_cc1 %s -emit-llvm -cl-mad-enable -o - | FileCheck %s -check-prefix=MAD +// RUN: %clang_cc1 %s -emit-llvm -cl-no-signed-zeros -o - | FileCheck %s -check-prefix=NOSIGNED float spscalardiv(float a, float b) { // CHECK: @spscalardiv( - // NORMAL: fdiv float + // NORMAL: fdiv float // FAST: fdiv fast float // FINITE: fdiv nnan ninf float - // UNSAFE: fdiv nnan float + // UNSAFE: fdiv nnan nsz float + // MAD: fdiv float + // NOSIGNED: fdiv nsz float return a / b; } // CHECK: attributes +// NORMAL: "less-precise-fpmad"="false" // NORMAL: "no-infs-fp-math"="false" // NORMAL: "no-nans-fp-math"="false" +// NORMAL: "no-signed-zeros-fp-math"="false" // NORMAL: "unsafe-fp-math"="false" +// FAST: "less-precise-fpmad"="true" // FAST: "no-infs-fp-math"="true" // FAST: "no-nans-fp-math"="true" +// FAST: "no-signed-zeros-fp-math"="true" // FAST: "unsafe-fp-math"="true" +// FINITE: "less-precise-fpmad"="false" // FINITE: "no-infs-fp-math"="true" // FINITE: "no-nans-fp-math"="true" +// FINITE: "no-signed-zeros-fp-math"="false" // FINITE: "unsafe-fp-math"="false" +// UNSAFE: "less-precise-fpmad"="true" // UNSAFE: "no-infs-fp-math"="false" // UNSAFE: "no-nans-fp-math"="true" +// UNSAFE: "no-signed-zeros-fp-math"="true" // UNSAFE: "unsafe-fp-math"="true" +// MAD: "less-precise-fpmad"="true" +// MAD: "no-infs-fp-math"="false" +// MAD: "no-nans-fp-math"="false" +// MAD: "no-signed-zeros-fp-math"="false" +// MAD: "unsafe-fp-math"="false" + +// NOSIGNED: "less-precise-fpmad"="false" +// NOSIGNED: "no-infs-fp-math"="false" +// NOSIGNED: "no-nans-fp-math"="false" +// NOSIGNED: "no-signed-zeros-fp-math"="true" +// NOSIGNED: "unsafe-fp-math"="false" diff --git a/test/CodeGenOpenCL/sampler.cl b/test/CodeGenOpenCL/sampler.cl index 6fc8c2c68446..3a7319cd78d2 100644 --- a/test/CodeGenOpenCL/sampler.cl +++ b/test/CodeGenOpenCL/sampler.cl @@ -54,4 +54,8 @@ kernel void foo(sampler_t smp_par) { 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]]) + + fnc4smp(5); + // CHECK: [[SAMP:%[0-9]+]] = call %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 5) + // CHECK: call spir_func void @fnc4smp(%opencl.sampler_t addrspace(2)* [[SAMP]]) } diff --git a/test/CodeGenOpenCL/vla.cl b/test/CodeGenOpenCL/vla.cl index cbf98445ce00..5d3599fc3cce 100644 --- a/test/CodeGenOpenCL/vla.cl +++ b/test/CodeGenOpenCL/vla.cl @@ -1,4 +1,6 @@ -// RUN: %clang_cc1 -emit-llvm -triple "spir-unknown-unknown" -O0 -cl-std=CL2.0 -o - %s | FileCheck %s +// RUN: %clang_cc1 -emit-llvm -triple "spir-unknown-unknown" -O0 -cl-std=CL2.0 -o - %s | FileCheck -check-prefixes=CHECK,SPIR %s +// RUN: %clang_cc1 -emit-llvm -triple amdgcn-amd-amdhsa-opencl -O0 -cl-std=CL2.0 -o - %s | FileCheck -check-prefixes=CHECK,SPIR %s +// RUN: %clang_cc1 -emit-llvm -triple amdgcn-amd-amdhsa-amdgizcl -O0 -cl-std=CL2.0 -o - %s | FileCheck -check-prefixes=CHECK,GIZ %s constant int sz0 = 5; // CHECK: @sz0 = addrspace(2) constant i32 5 @@ -11,8 +13,12 @@ const constant int sz2 = 8; kernel void testvla() { int vla0[sz0]; -// CHECK: %vla0 = alloca [5 x i32] +// SPIR: %vla0 = alloca [5 x i32] +// SPIR-NOT: %vla0 = alloca [5 x i32]{{.*}}addrspace +// GIZ: %vla0 = alloca [5 x i32]{{.*}}addrspace(5) char vla1[sz1]; -// CHECK: %vla1 = alloca [16 x i8] +// SPIR: %vla1 = alloca [16 x i8] +// SPIR-NOT: %vla1 = alloca [16 x i8]{{.*}}addrspace +// GIZ: %vla1 = alloca [16 x i8]{{.*}}addrspace(5) local short vla2[sz2]; } |