diff options
| author | Nikita Popov <npopov@redhat.com> | 2023-01-05 10:57:30 +0100 |
|---|---|---|
| committer | Nikita Popov <npopov@redhat.com> | 2023-01-05 10:57:30 +0100 |
| commit | aae20a7421c5393316c25a8b614b370859c1a18f (patch) | |
| tree | 33d90f3c9ef9b9093048f286b89b824e1ffdb89b | |
| parent | ccb6e0a51c15ca07928559d412a39093cf074366 (diff) | |
| download | llvm-aae20a7421c5393316c25a8b614b370859c1a18f.zip llvm-aae20a7421c5393316c25a8b614b370859c1a18f.tar.gz llvm-aae20a7421c5393316c25a8b614b370859c1a18f.tar.bz2 | |
[CodeGenOpenCL] Convert some tests to opaque pointers (NFC)
45 files changed, 483 insertions, 505 deletions
diff --git a/clang/test/CodeGenOpenCL/addr-space-struct-arg.cl b/clang/test/CodeGenOpenCL/addr-space-struct-arg.cl index 2f45dc2..a806c5f 100644 --- a/clang/test/CodeGenOpenCL/addr-space-struct-arg.cl +++ b/clang/test/CodeGenOpenCL/addr-space-struct-arg.cl @@ -1,9 +1,9 @@ -// RUN: %clang_cc1 -no-opaque-pointers %s -emit-llvm -o - -O0 -ffake-address-space-map -triple i686-pc-darwin | FileCheck -enable-var-scope -check-prefixes=ALL,X86 %s -// RUN: %clang_cc1 -no-opaque-pointers %s -emit-llvm -o - -O0 -triple amdgcn | FileCheck -enable-var-scope -check-prefixes=ALL,AMDGCN %s -// RUN: %clang_cc1 -no-opaque-pointers %s -emit-llvm -o - -cl-std=CL2.0 -O0 -triple amdgcn | FileCheck -enable-var-scope -check-prefixes=ALL,AMDGCN,AMDGCN20 %s -// RUN: %clang_cc1 -no-opaque-pointers %s -emit-llvm -o - -cl-std=CL1.2 -O0 -triple spir-unknown-unknown-unknown | FileCheck -enable-var-scope -check-prefixes=SPIR %s -// RUN: %clang_cc1 -no-opaque-pointers %s -emit-llvm -o - -cl-std=CL3.0 -O0 -triple amdgcn -cl-ext=+__opencl_c_program_scope_global_variables | FileCheck -enable-var-scope -check-prefixes=ALL,AMDGCN,AMDGCN20 %s -// RUN: %clang_cc1 -no-opaque-pointers %s -emit-llvm -o - -cl-std=CL3.0 -O0 -triple amdgcn | FileCheck -enable-var-scope -check-prefixes=ALL,AMDGCN %s +// RUN: %clang_cc1 %s -emit-llvm -o - -O0 -ffake-address-space-map -triple i686-pc-darwin | FileCheck -enable-var-scope -check-prefixes=ALL,X86 %s +// RUN: %clang_cc1 %s -emit-llvm -o - -O0 -triple amdgcn | FileCheck -enable-var-scope -check-prefixes=ALL,AMDGCN %s +// RUN: %clang_cc1 %s -emit-llvm -o - -cl-std=CL2.0 -O0 -triple amdgcn | FileCheck -enable-var-scope -check-prefixes=ALL,AMDGCN,AMDGCN20 %s +// RUN: %clang_cc1 %s -emit-llvm -o - -cl-std=CL1.2 -O0 -triple spir-unknown-unknown-unknown | FileCheck -enable-var-scope -check-prefixes=SPIR %s +// RUN: %clang_cc1 %s -emit-llvm -o - -cl-std=CL3.0 -O0 -triple amdgcn -cl-ext=+__opencl_c_program_scope_global_variables | FileCheck -enable-var-scope -check-prefixes=ALL,AMDGCN,AMDGCN20 %s +// RUN: %clang_cc1 %s -emit-llvm -o - -cl-std=CL3.0 -O0 -triple amdgcn | FileCheck -enable-var-scope -check-prefixes=ALL,AMDGCN %s typedef int int2 __attribute__((ext_vector_type(2))); @@ -45,7 +45,7 @@ struct LargeStructTwoMember { struct LargeStructOneMember g_s; #endif -// X86-LABEL: define{{.*}} void @foo(%struct.Mat4X4* noalias sret(%struct.Mat4X4) align 4 %agg.result, %struct.Mat3X3* noundef byval(%struct.Mat3X3) align 4 %in) +// X86-LABEL: define{{.*}} void @foo(ptr noalias sret(%struct.Mat4X4) align 4 %agg.result, ptr noundef byval(%struct.Mat3X3) align 4 %in) // AMDGCN-LABEL: define{{.*}} %struct.Mat4X4 @foo([9 x i32] %in.coerce) Mat4X4 __attribute__((noinline)) foo(Mat3X3 in) { Mat4X4 out; @@ -55,18 +55,18 @@ Mat4X4 __attribute__((noinline)) foo(Mat3X3 in) { // ALL-LABEL: define {{.*}} void @ker // Expect two mem copies: one for the argument "in", and one for // the return value. -// X86: call void @llvm.memcpy.p0i8.p1i8.i32(i8* -// X86: call void @llvm.memcpy.p1i8.p0i8.i32(i8 addrspace(1)* +// X86: call void @llvm.memcpy.p0.p1.i32(ptr +// X86: call void @llvm.memcpy.p1.p0.i32(ptr addrspace(1) -// AMDGCN: load [9 x i32], [9 x i32] addrspace(1)* +// AMDGCN: load [9 x i32], ptr addrspace(1) // AMDGCN: call %struct.Mat4X4 @foo([9 x i32] -// AMDGCN: call void @llvm.memcpy.p1i8.p5i8.i64(i8 addrspace(1)* +// AMDGCN: call void @llvm.memcpy.p1.p5.i64(ptr addrspace(1) kernel void ker(global Mat3X3 *in, global Mat4X4 *out) { out[0] = foo(in[1]); } -// X86-LABEL: define{{.*}} void @foo_large(%struct.Mat64X64* noalias sret(%struct.Mat64X64) align 4 %agg.result, %struct.Mat32X32* noundef byval(%struct.Mat32X32) align 4 %in) -// AMDGCN-LABEL: define{{.*}} void @foo_large(%struct.Mat64X64 addrspace(5)* noalias sret(%struct.Mat64X64) align 4 %agg.result, %struct.Mat32X32 addrspace(5)* noundef byval(%struct.Mat32X32) align 4 %in) +// X86-LABEL: define{{.*}} void @foo_large(ptr noalias sret(%struct.Mat64X64) align 4 %agg.result, ptr noundef byval(%struct.Mat32X32) align 4 %in) +// AMDGCN-LABEL: define{{.*}} void @foo_large(ptr addrspace(5) noalias sret(%struct.Mat64X64) align 4 %agg.result, ptr addrspace(5) noundef byval(%struct.Mat32X32) align 4 %in) Mat64X64 __attribute__((noinline)) foo_large(Mat32X32 in) { Mat64X64 out; return out; @@ -75,10 +75,10 @@ Mat64X64 __attribute__((noinline)) foo_large(Mat32X32 in) { // ALL-LABEL: define {{.*}} void @ker_large // Expect two mem copies: one for the argument "in", and one for // the return value. -// X86: call void @llvm.memcpy.p0i8.p1i8.i32(i8* -// X86: call void @llvm.memcpy.p1i8.p0i8.i32(i8 addrspace(1)* -// AMDGCN: call void @llvm.memcpy.p5i8.p1i8.i64(i8 addrspace(5)* -// AMDGCN: call void @llvm.memcpy.p1i8.p5i8.i64(i8 addrspace(1)* +// X86: call void @llvm.memcpy.p0.p1.i32(ptr +// X86: call void @llvm.memcpy.p1.p0.i32(ptr addrspace(1) +// AMDGCN: call void @llvm.memcpy.p5.p1.i64(ptr addrspace(5) +// AMDGCN: call void @llvm.memcpy.p1.p5.i64(ptr addrspace(1) kernel void ker_large(global Mat32X32 *in, global Mat64X64 *out) { out[0] = foo_large(in[1]); } @@ -88,18 +88,17 @@ void FuncOneMember(struct StructOneMember u) { u.x = (int2)(0, 0); } -// AMDGCN-LABEL: define{{.*}} void @FuncOneLargeMember(%struct.LargeStructOneMember addrspace(5)* noundef byval(%struct.LargeStructOneMember) align 8 %u) +// AMDGCN-LABEL: define{{.*}} void @FuncOneLargeMember(ptr addrspace(5) noundef byval(%struct.LargeStructOneMember) align 8 %u) // AMDGCN-NOT: addrspacecast -// AMDGCN: store <2 x i32> %{{.*}}, <2 x i32> addrspace(5)* +// AMDGCN: store <2 x i32> %{{.*}}, ptr addrspace(5) void FuncOneLargeMember(struct LargeStructOneMember u) { u.x[0] = (int2)(0, 0); } // AMDGCN20-LABEL: define{{.*}} void @test_indirect_arg_globl() // AMDGCN20: %[[byval_temp:.*]] = alloca %struct.LargeStructOneMember, align 8, addrspace(5) -// AMDGCN20: %[[r0:.*]] = bitcast %struct.LargeStructOneMember addrspace(5)* %[[byval_temp]] to i8 addrspace(5)* -// AMDGCN20: call void @llvm.memcpy.p5i8.p1i8.i64(i8 addrspace(5)* align 8 %[[r0]], i8 addrspace(1)* align 8 bitcast (%struct.LargeStructOneMember addrspace(1)* @g_s to i8 addrspace(1)*), i64 800, i1 false) -// AMDGCN20: call void @FuncOneLargeMember(%struct.LargeStructOneMember addrspace(5)* noundef byval(%struct.LargeStructOneMember) align 8 %[[byval_temp]]) +// AMDGCN20: call void @llvm.memcpy.p5.p1.i64(ptr addrspace(5) align 8 %[[byval_temp]], ptr addrspace(1) align 8 @g_s, i64 800, i1 false) +// AMDGCN20: call void @FuncOneLargeMember(ptr addrspace(5) noundef byval(%struct.LargeStructOneMember) align 8 %[[byval_temp]]) #if (__OPENCL_C_VERSION__ == 200) || (__OPENCL_C_VERSION__ >= 300 && defined(__opencl_c_program_scope_global_variables)) void test_indirect_arg_globl(void) { FuncOneLargeMember(g_s); @@ -108,9 +107,8 @@ void test_indirect_arg_globl(void) { // AMDGCN-LABEL: define{{.*}} amdgpu_kernel void @test_indirect_arg_local() // AMDGCN: %[[byval_temp:.*]] = alloca %struct.LargeStructOneMember, align 8, addrspace(5) -// AMDGCN: %[[r0:.*]] = bitcast %struct.LargeStructOneMember addrspace(5)* %[[byval_temp]] to i8 addrspace(5)* -// AMDGCN: call void @llvm.memcpy.p5i8.p3i8.i64(i8 addrspace(5)* align 8 %[[r0]], i8 addrspace(3)* align 8 bitcast (%struct.LargeStructOneMember addrspace(3)* @test_indirect_arg_local.l_s to i8 addrspace(3)*), i64 800, i1 false) -// AMDGCN: call void @FuncOneLargeMember(%struct.LargeStructOneMember addrspace(5)* noundef byval(%struct.LargeStructOneMember) align 8 %[[byval_temp]]) +// AMDGCN: call void @llvm.memcpy.p5.p3.i64(ptr addrspace(5) align 8 %[[byval_temp]], ptr addrspace(3) align 8 @test_indirect_arg_local.l_s, i64 800, i1 false) +// AMDGCN: call void @FuncOneLargeMember(ptr addrspace(5) noundef byval(%struct.LargeStructOneMember) align 8 %[[byval_temp]]) kernel void test_indirect_arg_local(void) { local struct LargeStructOneMember l_s; FuncOneLargeMember(l_s); @@ -119,7 +117,7 @@ kernel void test_indirect_arg_local(void) { // AMDGCN-LABEL: define{{.*}} void @test_indirect_arg_private() // AMDGCN: %[[p_s:.*]] = alloca %struct.LargeStructOneMember, align 8, addrspace(5) // AMDGCN-NOT: @llvm.memcpy -// AMDGCN-NEXT: call void @FuncOneLargeMember(%struct.LargeStructOneMember addrspace(5)* noundef byval(%struct.LargeStructOneMember) align 8 %[[p_s]]) +// AMDGCN-NEXT: call void @FuncOneLargeMember(ptr addrspace(5) noundef byval(%struct.LargeStructOneMember) align 8 %[[p_s]]) void test_indirect_arg_private(void) { struct LargeStructOneMember p_s; FuncOneLargeMember(p_s); @@ -128,14 +126,14 @@ void test_indirect_arg_private(void) { // AMDGCN-LABEL: define{{.*}} amdgpu_kernel void @KernelOneMember // AMDGCN-SAME: (<2 x i32> %[[u_coerce:.*]]) // AMDGCN: %[[u:.*]] = alloca %struct.StructOneMember, align 8, addrspace(5) -// AMDGCN: %[[coerce_dive:.*]] = getelementptr inbounds %struct.StructOneMember, %struct.StructOneMember addrspace(5)* %[[u]], i32 0, i32 0 -// AMDGCN: store <2 x i32> %[[u_coerce]], <2 x i32> addrspace(5)* %[[coerce_dive]] +// AMDGCN: %[[coerce_dive:.*]] = getelementptr inbounds %struct.StructOneMember, ptr addrspace(5) %[[u]], i32 0, i32 0 +// AMDGCN: store <2 x i32> %[[u_coerce]], ptr addrspace(5) %[[coerce_dive]] // AMDGCN: call void @FuncOneMember(<2 x i32> kernel void KernelOneMember(struct StructOneMember u) { FuncOneMember(u); } -// SPIR: call void @llvm.memcpy.p0i8.p1i8.i32 +// SPIR: call void @llvm.memcpy.p0.p1.i32 // SPIR-NOT: addrspacecast kernel void KernelOneMemberSpir(global struct StructOneMember* u) { FuncOneMember(*u); @@ -143,8 +141,8 @@ kernel void KernelOneMemberSpir(global struct StructOneMember* u) { // AMDGCN-LABEL: define{{.*}} amdgpu_kernel void @KernelLargeOneMember( // AMDGCN: %[[U:.*]] = alloca %struct.LargeStructOneMember, align 8, addrspace(5) -// AMDGCN: store %struct.LargeStructOneMember %u.coerce, %struct.LargeStructOneMember addrspace(5)* %[[U]], align 8 -// AMDGCN: call void @FuncOneLargeMember(%struct.LargeStructOneMember addrspace(5)* noundef byval(%struct.LargeStructOneMember) align 8 %[[U]]) +// AMDGCN: store %struct.LargeStructOneMember %u.coerce, ptr addrspace(5) %[[U]], align 8 +// AMDGCN: call void @FuncOneLargeMember(ptr addrspace(5) noundef byval(%struct.LargeStructOneMember) align 8 %[[U]]) kernel void KernelLargeOneMember(struct LargeStructOneMember u) { FuncOneLargeMember(u); } @@ -154,7 +152,7 @@ void FuncTwoMember(struct StructTwoMember u) { u.y = (int2)(0, 0); } -// AMDGCN-LABEL: define{{.*}} void @FuncLargeTwoMember(%struct.LargeStructTwoMember addrspace(5)* noundef byval(%struct.LargeStructTwoMember) align 8 %u) +// AMDGCN-LABEL: define{{.*}} void @FuncLargeTwoMember(ptr addrspace(5) noundef byval(%struct.LargeStructTwoMember) align 8 %u) void FuncLargeTwoMember(struct LargeStructTwoMember u) { u.y[0] = (int2)(0, 0); } @@ -162,8 +160,8 @@ void FuncLargeTwoMember(struct LargeStructTwoMember u) { // AMDGCN-LABEL: define{{.*}} amdgpu_kernel void @KernelTwoMember // AMDGCN-SAME: (%struct.StructTwoMember %[[u_coerce:.*]]) // AMDGCN: %[[u:.*]] = alloca %struct.StructTwoMember, align 8, addrspace(5) -// AMDGCN: %[[LD0:.*]] = load <2 x i32>, <2 x i32> addrspace(5)* -// AMDGCN: %[[LD1:.*]] = load <2 x i32>, <2 x i32> addrspace(5)* +// AMDGCN: %[[LD0:.*]] = load <2 x i32>, ptr addrspace(5) +// AMDGCN: %[[LD1:.*]] = load <2 x i32>, ptr addrspace(5) // AMDGCN: call void @FuncTwoMember(<2 x i32> %[[LD0]], <2 x i32> %[[LD1]]) kernel void KernelTwoMember(struct StructTwoMember u) { FuncTwoMember(u); @@ -172,8 +170,8 @@ kernel void KernelTwoMember(struct StructTwoMember u) { // AMDGCN-LABEL: define{{.*}} amdgpu_kernel void @KernelLargeTwoMember // AMDGCN-SAME: (%struct.LargeStructTwoMember %[[u_coerce:.*]]) // AMDGCN: %[[u:.*]] = alloca %struct.LargeStructTwoMember, align 8, addrspace(5) -// AMDGCN: store %struct.LargeStructTwoMember %[[u_coerce]], %struct.LargeStructTwoMember addrspace(5)* %[[u]] -// AMDGCN: call void @FuncLargeTwoMember(%struct.LargeStructTwoMember addrspace(5)* noundef byval(%struct.LargeStructTwoMember) align 8 %[[u]]) +// AMDGCN: store %struct.LargeStructTwoMember %[[u_coerce]], ptr addrspace(5) %[[u]] +// AMDGCN: call void @FuncLargeTwoMember(ptr addrspace(5) noundef byval(%struct.LargeStructTwoMember) align 8 %[[u]]) kernel void KernelLargeTwoMember(struct LargeStructTwoMember u) { FuncLargeTwoMember(u); } diff --git a/clang/test/CodeGenOpenCL/address-space-constant-initializers.cl b/clang/test/CodeGenOpenCL/address-space-constant-initializers.cl index 8b488d2..14fbcba 100644 --- a/clang/test/CodeGenOpenCL/address-space-constant-initializers.cl +++ b/clang/test/CodeGenOpenCL/address-space-constant-initializers.cl @@ -1,5 +1,5 @@ -// RUN: %clang_cc1 -no-opaque-pointers %s -ffake-address-space-map -emit-llvm -o - | FileCheck -check-prefix=FAKE %s -// RUN: %clang_cc1 -no-opaque-pointers %s -triple amdgcn-amd-amdhsa -emit-llvm -o - | FileCheck -check-prefix=AMDGCN %s +// RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck -check-prefix=FAKE %s +// RUN: %clang_cc1 %s -triple amdgcn-amd-amdhsa -emit-llvm -o - | FileCheck -check-prefix=AMDGCN %s typedef struct { int i; @@ -12,10 +12,10 @@ typedef struct { __constant float* constant_float_ptr; } ConstantArrayPointerStruct; -// FAKE: %struct.ConstantArrayPointerStruct = type { float addrspace(2)* } -// FAKE: addrspace(2) constant %struct.ConstantArrayPointerStruct { float addrspace(2)* bitcast (i8 addrspace(2)* getelementptr (i8, i8 addrspace(2)* bitcast (%struct.ArrayStruct addrspace(2)* @constant_array_struct to i8 addrspace(2)*), i64 4) to float addrspace(2)*) } -// AMDGCN: %struct.ConstantArrayPointerStruct = type { float addrspace(4)* } -// AMDGCN: addrspace(4) constant %struct.ConstantArrayPointerStruct { float addrspace(4)* bitcast (i8 addrspace(4)* getelementptr (i8, i8 addrspace(4)* bitcast (%struct.ArrayStruct addrspace(4)* @constant_array_struct to i8 addrspace(4)*), i64 4) to float addrspace(4)*) } +// FAKE: %struct.ConstantArrayPointerStruct = type { ptr addrspace(2) } +// FAKE: addrspace(2) constant %struct.ConstantArrayPointerStruct { ptr addrspace(2) getelementptr (i8, ptr addrspace(2) @constant_array_struct, i64 4) } +// AMDGCN: %struct.ConstantArrayPointerStruct = type { ptr addrspace(4) } +// AMDGCN: addrspace(4) constant %struct.ConstantArrayPointerStruct { ptr addrspace(4) getelementptr (i8, ptr addrspace(4) @constant_array_struct, i64 4) } // Bug 18567 __constant ConstantArrayPointerStruct constant_array_pointer_struct = { &constant_array_struct.f diff --git a/clang/test/CodeGenOpenCL/address-spaces.cl b/clang/test/CodeGenOpenCL/address-spaces.cl index 7ad8220..5b2a95c 100644 --- a/clang/test/CodeGenOpenCL/address-spaces.cl +++ b/clang/test/CodeGenOpenCL/address-spaces.cl @@ -1,17 +1,17 @@ -// RUN: %clang_cc1 -no-opaque-pointers %s -O0 -ffake-address-space-map -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,SPIR -// RUN: %clang_cc1 -no-opaque-pointers %s -O0 -cl-std=CL3.0 -cl-ext=-all -ffake-address-space-map -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,SPIR -// RUN: %clang_cc1 -no-opaque-pointers %s -O0 -cl-std=clc++2021 -cl-ext=-all -ffake-address-space-map -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,SPIR -// RUN: %clang_cc1 -no-opaque-pointers %s -O0 -DCL20 -cl-std=CL2.0 -ffake-address-space-map -emit-llvm -o - | FileCheck %s --check-prefixes=CL20,CL20SPIR -// RUN: %clang_cc1 -no-opaque-pointers %s -O0 -triple amdgcn-amd-amdhsa -emit-llvm -o - | FileCheck --check-prefixes=CHECK,AMDGCN %s -// RUN: %clang_cc1 -no-opaque-pointers %s -O0 -triple amdgcn-amd-amdhsa -cl-std=CL3.0 -emit-llvm -o - | FileCheck --check-prefixes=CHECK,AMDGCN %s -// RUN: %clang_cc1 -no-opaque-pointers %s -O0 -triple amdgcn-amd-amdhsa -DCL20 -cl-std=CL2.0 -emit-llvm -o - | FileCheck %s --check-prefixes=CL20,CL20AMDGCN -// RUN: %clang_cc1 -no-opaque-pointers %s -O0 -triple amdgcn-mesa-mesa3d -emit-llvm -o - | FileCheck --check-prefixes=CHECK,AMDGCN %s -// RUN: %clang_cc1 -no-opaque-pointers %s -O0 -triple amdgcn-mesa-mesa3d -cl-std=CL3.0 -emit-llvm -o - | FileCheck --check-prefixes=CHECK,AMDGCN %s -// RUN: %clang_cc1 -no-opaque-pointers %s -O0 -triple r600-- -emit-llvm -o - | FileCheck --check-prefixes=CHECK,AMDGCN %s -// RUN: %clang_cc1 -no-opaque-pointers %s -O0 -triple r600-- -emit-llvm -cl-std=CL3.0 -o - | FileCheck --check-prefixes=CHECK,AMDGCN %s - -// SPIR: %struct.S = type { i32, i32, i32* } -// CL20SPIR: %struct.S = type { i32, i32, i32 addrspace(4)* } +// RUN: %clang_cc1 %s -O0 -ffake-address-space-map -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,SPIR +// RUN: %clang_cc1 %s -O0 -cl-std=CL3.0 -cl-ext=-all -ffake-address-space-map -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,SPIR +// RUN: %clang_cc1 %s -O0 -cl-std=clc++2021 -cl-ext=-all -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 -emit-llvm -o - | FileCheck --check-prefixes=CHECK,AMDGCN %s +// RUN: %clang_cc1 %s -O0 -triple amdgcn-amd-amdhsa -cl-std=CL3.0 -emit-llvm -o - | FileCheck --check-prefixes=CHECK,AMDGCN %s +// RUN: %clang_cc1 %s -O0 -triple amdgcn-amd-amdhsa -DCL20 -cl-std=CL2.0 -emit-llvm -o - | FileCheck %s --check-prefixes=CL20,CL20AMDGCN +// RUN: %clang_cc1 %s -O0 -triple amdgcn-mesa-mesa3d -emit-llvm -o - | FileCheck --check-prefixes=CHECK,AMDGCN %s +// RUN: %clang_cc1 %s -O0 -triple amdgcn-mesa-mesa3d -cl-std=CL3.0 -emit-llvm -o - | FileCheck --check-prefixes=CHECK,AMDGCN %s +// RUN: %clang_cc1 %s -O0 -triple r600-- -emit-llvm -o - | FileCheck --check-prefixes=CHECK,AMDGCN %s +// RUN: %clang_cc1 %s -O0 -triple r600-- -emit-llvm -cl-std=CL3.0 -o - | FileCheck --check-prefixes=CHECK,AMDGCN %s + +// SPIR: %struct.S = type { i32, i32, ptr } +// CL20SPIR: %struct.S = type { i32, i32, ptr addrspace(4) } struct S { int x; int y; @@ -28,54 +28,54 @@ struct S { struct S g_s; #endif -// SPIR: i32* %arg -// AMDGCN: i32 addrspace(5)* %arg +// SPIR: ptr %arg +// AMDGCN: ptr addrspace(5) %arg void f__p(__private int *arg) {} -// CHECK: i32 addrspace(1)* %arg +// CHECK: ptr addrspace(1) %arg void f__g(__global int *arg) {} -// CHECK: i32 addrspace(3)* %arg +// CHECK: ptr addrspace(3) %arg void f__l(__local int *arg) {} -// SPIR: i32 addrspace(2)* %arg -// AMDGCN: i32 addrspace(4)* %arg +// SPIR: ptr addrspace(2) %arg +// AMDGCN: ptr addrspace(4) %arg void f__c(__constant int *arg) {} -// SPIR: i32* %arg -// AMDGCN: i32 addrspace(5)* %arg +// SPIR: ptr %arg +// AMDGCN: ptr addrspace(5) %arg void fp(private int *arg) {} -// CHECK: i32 addrspace(1)* %arg +// CHECK: ptr addrspace(1) %arg void fg(global int *arg) {} -// CHECK: i32 addrspace(3)* %arg +// CHECK: ptr addrspace(3) %arg void fl(local int *arg) {} -// SPIR: i32 addrspace(2)* %arg -// AMDGCN: i32 addrspace(4)* %arg +// SPIR: ptr addrspace(2) %arg +// AMDGCN: ptr addrspace(4) %arg void fc(constant int *arg) {} -// SPIR: i32 addrspace(5)* %arg -// AMDGCN: i32 addrspace(1)* %arg +// SPIR: ptr addrspace(5) %arg +// AMDGCN: ptr addrspace(1) %arg void fd(__attribute__((opencl_global_device)) int *arg) {} -// SPIR: i32 addrspace(6)* %arg -// AMDGCN: i32 addrspace(1)* %arg +// SPIR: ptr addrspace(6) %arg +// AMDGCN: ptr addrspace(1) %arg void fh(__attribute__((opencl_global_host)) int *arg) {} #ifdef CL20 int i; // CL20-DAG: @i = {{(dso_local )?}}addrspace(1) global i32 0 int *ptr; -// CL20SPIR-DAG: @ptr = {{(common )?}}{{(dso_local )?}}addrspace(1) global i32 addrspace(4)* null -// CL20AMDGCN-DAG: @ptr = {{(dso_local )?}}addrspace(1) global i32* null +// CL20SPIR-DAG: @ptr = {{(common )?}}{{(dso_local )?}}addrspace(1) global ptr addrspace(4) null +// CL20AMDGCN-DAG: @ptr = {{(dso_local )?}}addrspace(1) global ptr null #endif -// SPIR: i32* noundef %arg -// AMDGCN: i32 addrspace(5)* noundef %arg -// CL20SPIR-DAG: i32 addrspace(4)* noundef %arg -// CL20AMDGCN-DAG: i32* noundef %arg +// SPIR: ptr noundef %arg +// AMDGCN: ptr addrspace(5) noundef %arg +// CL20SPIR-DAG: ptr addrspace(4) noundef %arg +// CL20AMDGCN-DAG: ptr noundef %arg void f(int *arg) { int i; @@ -92,7 +92,7 @@ void f(int *arg) { typedef int int_td; typedef int *intp_td; -// SPIR: define {{(dso_local )?}}void @{{.*}}test_typedef{{.*}}(i32 addrspace(1)* noundef %x, i32 addrspace(2)* noundef %y, i32* noundef %z) +// SPIR: define {{(dso_local )?}}void @{{.*}}test_typedef{{.*}}(ptr addrspace(1) noundef %x, ptr addrspace(2) noundef %y, ptr noundef %z) void test_typedef(global int_td *x, constant int_td *y, intp_td z) { *x = *y; *z = 0; @@ -100,14 +100,14 @@ void test_typedef(global int_td *x, constant int_td *y, intp_td z) { // SPIR: define {{(dso_local )?}}void @{{.*}}test_struct{{.*}}() void test_struct() { - // SPIR: %ps = alloca %struct.S* - // CL20SPIR: %ps = alloca %struct.S addrspace(4)* + // SPIR: %ps = alloca ptr + // CL20SPIR: %ps = alloca ptr addrspace(4) struct S *ps; - // SPIR: store i32 0, i32* %x - // CL20SPIR: store i32 0, i32 addrspace(4)* %x + // SPIR: store i32 0, ptr %x + // CL20SPIR: store i32 0, ptr addrspace(4) %x ps->x = 0; #ifdef CL20 - // CL20SPIR: store i32 0, i32 addrspace(1)* getelementptr inbounds (%struct.S, %struct.S addrspace(1)* @g_s, i32 0, i32 0) + // CL20SPIR: store i32 0, ptr addrspace(1) @g_s g_s.x = 0; #endif } diff --git a/clang/test/CodeGenOpenCL/amdgcn-automatic-variable.cl b/clang/test/CodeGenOpenCL/amdgcn-automatic-variable.cl index 88754c7..f26495bc 100644 --- a/clang/test/CodeGenOpenCL/amdgcn-automatic-variable.cl +++ b/clang/test/CodeGenOpenCL/amdgcn-automatic-variable.cl @@ -1,17 +1,17 @@ -// RUN: %clang_cc1 -no-opaque-pointers -O0 -cl-std=CL1.2 -triple amdgcn---amdgizcl -emit-llvm %s -o - | FileCheck -check-prefixes=CHECK,CL12 %s -// RUN: %clang_cc1 -no-opaque-pointers -O0 -cl-std=CL2.0 -triple amdgcn---amdgizcl -emit-llvm %s -o - | FileCheck -check-prefixes=CHECK,CL20 %s +// RUN: %clang_cc1 -O0 -cl-std=CL1.2 -triple amdgcn---amdgizcl -emit-llvm %s -o - | FileCheck -check-prefixes=CHECK,CL12 %s +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn---amdgizcl -emit-llvm %s -o - | FileCheck -check-prefixes=CHECK,CL20 %s -// CL12-LABEL: define{{.*}} void @func1(i32 addrspace(5)* noundef %x) -// CL20-LABEL: define{{.*}} void @func1(i32* noundef %x) +// CL12-LABEL: define{{.*}} void @func1(ptr addrspace(5) noundef %x) +// CL20-LABEL: define{{.*}} void @func1(ptr noundef %x) void func1(int *x) { - // CL12: %[[x_addr:.*]] = alloca i32 addrspace(5)*{{.*}}addrspace(5) - // CL12: store i32 addrspace(5)* %x, i32 addrspace(5)* addrspace(5)* %[[x_addr]] - // CL12: %[[r0:.*]] = load i32 addrspace(5)*, i32 addrspace(5)* addrspace(5)* %[[x_addr]] - // CL12: store i32 1, i32 addrspace(5)* %[[r0]] - // CL20: %[[x_addr:.*]] = alloca i32*{{.*}}addrspace(5) - // CL20: store i32* %x, i32* addrspace(5)* %[[x_addr]] - // CL20: %[[r0:.*]] = load i32*, i32* addrspace(5)* %[[x_addr]] - // CL20: store i32 1, i32* %[[r0]] + // CL12: %[[x_addr:.*]] = alloca ptr addrspace(5){{.*}}addrspace(5) + // CL12: store ptr addrspace(5) %x, ptr addrspace(5) %[[x_addr]] + // CL12: %[[r0:.*]] = load ptr addrspace(5), ptr addrspace(5) %[[x_addr]] + // CL12: store i32 1, ptr addrspace(5) %[[r0]] + // CL20: %[[x_addr:.*]] = alloca ptr{{.*}}addrspace(5) + // CL20: store ptr %x, ptr addrspace(5) %[[x_addr]] + // CL20: %[[r0:.*]] = load ptr, ptr addrspace(5) %[[x_addr]] + // CL20: store i32 1, ptr %[[r0]] *x = 1; } @@ -20,49 +20,48 @@ void func2(void) { // CHECK: %lv1 = alloca i32, align 4, addrspace(5) // CHECK: %lv2 = alloca i32, align 4, addrspace(5) // CHECK: %la = alloca [100 x i32], align 4, addrspace(5) - // CL12: %lp1 = alloca i32 addrspace(5)*, align 4, addrspace(5) - // CL12: %lp2 = alloca i32 addrspace(5)*, align 4, addrspace(5) - // CL20: %lp1 = alloca i32*, align 8, addrspace(5) - // CL20: %lp2 = alloca i32*, align 8, addrspace(5) + // CL12: %lp1 = alloca ptr addrspace(5), align 4, addrspace(5) + // CL12: %lp2 = alloca ptr addrspace(5), align 4, addrspace(5) + // CL20: %lp1 = alloca ptr, align 8, addrspace(5) + // CL20: %lp2 = alloca ptr, align 8, addrspace(5) // CHECK: %lvc = alloca i32, align 4, addrspace(5) - // CHECK: store i32 1, i32 addrspace(5)* %lv1 + // CHECK: store i32 1, ptr addrspace(5) %lv1 int lv1; lv1 = 1; - // CHECK: store i32 2, i32 addrspace(5)* %lv2 + // CHECK: store i32 2, ptr addrspace(5) %lv2 int lv2 = 2; - // CHECK: %[[arrayidx:.*]] = getelementptr inbounds [100 x i32], [100 x i32] addrspace(5)* %la, i64 0, i64 0 - // CHECK: store i32 3, i32 addrspace(5)* %[[arrayidx]], align 4 + // CHECK: %[[arrayidx:.*]] = getelementptr inbounds [100 x i32], ptr addrspace(5) %la, i64 0, i64 0 + // CHECK: store i32 3, ptr addrspace(5) %[[arrayidx]], align 4 int la[100]; la[0] = 3; - // CL12: store i32 addrspace(5)* %lv1, i32 addrspace(5)* addrspace(5)* %lp1, align 4 - // CL20: %[[r0:.*]] = addrspacecast i32 addrspace(5)* %lv1 to i32* - // CL20: store i32* %[[r0]], i32* addrspace(5)* %lp1, align 8 + // CL12: store ptr addrspace(5) %lv1, ptr addrspace(5) %lp1, align 4 + // CL20: %[[r0:.*]] = addrspacecast ptr addrspace(5) %lv1 to ptr + // CL20: store ptr %[[r0]], ptr addrspace(5) %lp1, align 8 int *lp1 = &lv1; - // CHECK: %[[arraydecay:.*]] = getelementptr inbounds [100 x i32], [100 x i32] addrspace(5)* %la, i64 0, i64 0 - // CL12: store i32 addrspace(5)* %[[arraydecay]], i32 addrspace(5)* addrspace(5)* %lp2, align 4 - // CL20: %[[r1:.*]] = addrspacecast i32 addrspace(5)* %[[arraydecay]] to i32* - // CL20: store i32* %[[r1]], i32* addrspace(5)* %lp2, align 8 + // CHECK: %[[arraydecay:.*]] = getelementptr inbounds [100 x i32], ptr addrspace(5) %la, i64 0, i64 0 + // CL12: store ptr addrspace(5) %[[arraydecay]], ptr addrspace(5) %lp2, align 4 + // CL20: %[[r1:.*]] = addrspacecast ptr addrspace(5) %[[arraydecay]] to ptr + // CL20: store ptr %[[r1]], ptr addrspace(5) %lp2, align 8 int *lp2 = la; - // CL12: call void @func1(i32 addrspace(5)* noundef %lv1) - // CL20: %[[r2:.*]] = addrspacecast i32 addrspace(5)* %lv1 to i32* - // CL20: call void @func1(i32* noundef %[[r2]]) + // CL12: call void @func1(ptr addrspace(5) noundef %lv1) + // CL20: %[[r2:.*]] = addrspacecast ptr addrspace(5) %lv1 to ptr + // CL20: call void @func1(ptr noundef %[[r2]]) func1(&lv1); - // CHECK: store i32 4, i32 addrspace(5)* %lvc - // CHECK: store i32 4, i32 addrspace(5)* %lv1 + // CHECK: store i32 4, ptr addrspace(5) %lvc + // CHECK: store i32 4, ptr addrspace(5) %lv1 const int lvc = 4; lv1 = lvc; } // CHECK-LABEL: define{{.*}} void @func3() // CHECK: %a = alloca [16 x [1 x float]], align 4, addrspace(5) -// CHECK: %[[CAST:.+]] = bitcast [16 x [1 x float]] addrspace(5)* %a to i8 addrspace(5)* -// CHECK: call void @llvm.memset.p5i8.i64(i8 addrspace(5)* align 4 %[[CAST]], i8 0, i64 64, i1 false) +// CHECK: call void @llvm.memset.p5.i64(ptr addrspace(5) align 4 %a, i8 0, i64 64, i1 false) void func3(void) { float a[16][1] = {{0.}}; } diff --git a/clang/test/CodeGenOpenCL/amdgcn-non-temporal-store.cl b/clang/test/CodeGenOpenCL/amdgcn-non-temporal-store.cl index 72f1a00..3493860 100644 --- a/clang/test/CodeGenOpenCL/amdgcn-non-temporal-store.cl +++ b/clang/test/CodeGenOpenCL/amdgcn-non-temporal-store.cl @@ -1,7 +1,7 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s // CHECK-LABEL: @test_non_temporal_store_kernel -// CHECK: store i32 0, i32 addrspace(1)* %{{.*}}, align 4, !tbaa !{{.*}}, !nontemporal {{.*}} +// CHECK: store i32 0, ptr addrspace(1) %{{.*}}, align 4, !tbaa !{{.*}}, !nontemporal {{.*}} kernel void test_non_temporal_store_kernel(global unsigned int* io) { __builtin_nontemporal_store(0, io); diff --git a/clang/test/CodeGenOpenCL/amdgpu-call-kernel.cl b/clang/test/CodeGenOpenCL/amdgpu-call-kernel.cl index 22e0fd1c..8ad0bea 100755 --- a/clang/test/CodeGenOpenCL/amdgpu-call-kernel.cl +++ b/clang/test/CodeGenOpenCL/amdgpu-call-kernel.cl @@ -1,7 +1,7 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s -// CHECK: define{{.*}} amdgpu_kernel void @test_call_kernel(i32 addrspace(1)* nocapture noundef writeonly align 4 %out) -// CHECK: store i32 4, i32 addrspace(1)* %out, align 4 +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s +// CHECK: define{{.*}} amdgpu_kernel void @test_call_kernel(ptr addrspace(1) nocapture noundef writeonly align 4 %out) +// CHECK: store i32 4, ptr addrspace(1) %out, align 4 kernel void test_kernel(global int *out) { diff --git a/clang/test/CodeGenOpenCL/amdgpu-debug-info-variable-expression.cl b/clang/test/CodeGenOpenCL/amdgpu-debug-info-variable-expression.cl index 425f764..ba64544 100644 --- a/clang/test/CodeGenOpenCL/amdgpu-debug-info-variable-expression.cl +++ b/clang/test/CodeGenOpenCL/amdgpu-debug-info-variable-expression.cl @@ -1,6 +1,6 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang -Xclang -no-opaque-pointers -cl-std=CL2.0 -emit-llvm -g -O0 -S -nogpulib -target amdgcn-amd-amdhsa -mcpu=fiji -o - %s | FileCheck %s -// RUN: %clang -Xclang -no-opaque-pointers -cl-std=CL2.0 -emit-llvm -g -O0 -S -nogpulib -target amdgcn-amd-amdhsa-opencl -mcpu=fiji -o - %s | FileCheck %s +// RUN: %clang -cl-std=CL2.0 -emit-llvm -g -O0 -S -nogpulib -target amdgcn-amd-amdhsa -mcpu=fiji -o - %s | FileCheck %s +// RUN: %clang -cl-std=CL2.0 -emit-llvm -g -O0 -S -nogpulib -target amdgcn-amd-amdhsa-opencl -mcpu=fiji -o - %s | FileCheck %s // 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]], expr: !DIExpression()) @@ -52,31 +52,31 @@ 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)* addrspace(5)* {{.*}}, metadata ![[KERNELARG0]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}} + // CHECK-DAG: call void @llvm.dbg.declare(metadata ptr addrspace(5) {{.*}}, metadata ![[KERNELARG0]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !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(4)* addrspace(5)* {{.*}}, metadata ![[KERNELARG1]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}} + // CHECK-DAG: call void @llvm.dbg.declare(metadata ptr addrspace(5) {{.*}}, metadata ![[KERNELARG1]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !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)* addrspace(5)* {{.*}}, metadata ![[KERNELARG2]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}} + // CHECK-DAG: call void @llvm.dbg.declare(metadata ptr addrspace(5) {{.*}}, metadata ![[KERNELARG2]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !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)* addrspace(5)* {{.*}}, metadata ![[FUNCVAR0]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}} + // CHECK-DAG: call void @llvm.dbg.declare(metadata ptr addrspace(5) {{.*}}, metadata ![[FUNCVAR0]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !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(4)* addrspace(5)* {{.*}}, metadata ![[FUNCVAR1]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}} + // CHECK-DAG: call void @llvm.dbg.declare(metadata ptr addrspace(5) {{.*}}, metadata ![[FUNCVAR1]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !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)* addrspace(5)* {{.*}}, metadata ![[FUNCVAR2]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}} + // CHECK-DAG: call void @llvm.dbg.declare(metadata ptr addrspace(5) {{.*}}, metadata ![[FUNCVAR2]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !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 addrspace(5)* addrspace(5)* {{.*}}, metadata ![[FUNCVAR3]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}} + // CHECK-DAG: call void @llvm.dbg.declare(metadata ptr addrspace(5) {{.*}}, metadata ![[FUNCVAR3]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !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(5)* {{.*}}, metadata ![[FUNCVAR4]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}} + // CHECK-DAG: call void @llvm.dbg.declare(metadata ptr addrspace(5) {{.*}}, metadata ![[FUNCVAR4]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}} int *FuncVar4 = Tmp1; // CHECK-DAG: ![[FUNCVAR5:[0-9]+]] = distinct !DIGlobalVariable(name: "FuncVar5", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: true, isDefinition: true) @@ -112,18 +112,18 @@ kernel void kernel1( 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)* addrspace(5)* {{.*}}, metadata ![[FUNCVAR15]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}} + // CHECK-DAG: call void @llvm.dbg.declare(metadata ptr addrspace(5) {{.*}}, metadata ![[FUNCVAR15]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !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(4)* addrspace(5)* {{.*}}, metadata ![[FUNCVAR16]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}} + // CHECK-DAG: call void @llvm.dbg.declare(metadata ptr addrspace(5) {{.*}}, metadata ![[FUNCVAR16]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !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)* addrspace(5)* {{.*}}, metadata ![[FUNCVAR17]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}} + // CHECK-DAG: call void @llvm.dbg.declare(metadata ptr addrspace(5) {{.*}}, metadata ![[FUNCVAR17]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !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 addrspace(5)* addrspace(5)* {{.*}}, metadata ![[FUNCVAR18]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}} + // CHECK-DAG: call void @llvm.dbg.declare(metadata ptr addrspace(5) {{.*}}, metadata ![[FUNCVAR18]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !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(5)* {{.*}}, metadata ![[FUNCVAR19]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}} + // CHECK-DAG: call void @llvm.dbg.declare(metadata ptr addrspace(5) {{.*}}, metadata ![[FUNCVAR19]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}} int *private FuncVar19 = Tmp1; } diff --git a/clang/test/CodeGenOpenCL/amdgpu-printf.cl b/clang/test/CodeGenOpenCL/amdgpu-printf.cl index 1bc0e72..edf6dbf 100644 --- a/clang/test/CodeGenOpenCL/amdgpu-printf.cl +++ b/clang/test/CodeGenOpenCL/amdgpu-printf.cl @@ -1,11 +1,11 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py -// RUN: %clang_cc1 -no-opaque-pointers -cl-std=CL1.2 -triple amdgcn-amd-amdhsa -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL1.2 -triple amdgcn-amd-amdhsa -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s int printf(__constant const char* st, ...) __attribute__((format(printf, 1, 2))); // CHECK-LABEL: @test_printf_noargs( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL:%.*]] = call i32 (i8 addrspace(4)*, ...) @printf(i8 addrspace(4)* noundef getelementptr inbounds ([1 x i8], [1 x i8] addrspace(4)* @.str, i64 0, i64 0)) #[[ATTR4:[0-9]+]] +// CHECK-NEXT: [[CALL:%.*]] = call i32 (ptr addrspace(4), ...) @printf(ptr addrspace(4) noundef @.str) #[[ATTR4:[0-9]+]] // CHECK-NEXT: ret void // __kernel void test_printf_noargs() { @@ -15,9 +15,9 @@ __kernel void test_printf_noargs() { // CHECK-LABEL: @test_printf_int( // CHECK-NEXT: entry: // CHECK-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4, addrspace(5) -// CHECK-NEXT: store i32 [[I:%.*]], i32 addrspace(5)* [[I_ADDR]], align 4, !tbaa [[TBAA8:![0-9]+]] -// CHECK-NEXT: [[TMP0:%.*]] = load i32, i32 addrspace(5)* [[I_ADDR]], align 4, !tbaa [[TBAA8]] -// CHECK-NEXT: [[CALL:%.*]] = call i32 (i8 addrspace(4)*, ...) @printf(i8 addrspace(4)* noundef getelementptr inbounds ([3 x i8], [3 x i8] addrspace(4)* @.str.1, i64 0, i64 0), i32 noundef [[TMP0]]) #[[ATTR4]] +// CHECK-NEXT: store i32 [[I:%.*]], ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA8:![0-9]+]] +// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA8]] +// CHECK-NEXT: [[CALL:%.*]] = call i32 (ptr addrspace(4), ...) @printf(ptr addrspace(4) noundef @.str.1, i32 noundef [[TMP0]]) #[[ATTR4]] // CHECK-NEXT: ret void // __kernel void test_printf_int(int i) { @@ -28,16 +28,13 @@ __kernel void test_printf_int(int i) { // CHECK-NEXT: entry: // CHECK-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4, addrspace(5) // CHECK-NEXT: [[S:%.*]] = alloca [4 x i8], align 1, addrspace(5) -// CHECK-NEXT: store i32 [[I:%.*]], i32 addrspace(5)* [[I_ADDR]], align 4, !tbaa [[TBAA8]] -// CHECK-NEXT: [[TMP0:%.*]] = bitcast [4 x i8] addrspace(5)* [[S]] to i8 addrspace(5)* -// CHECK-NEXT: call void @llvm.lifetime.start.p5i8(i64 4, i8 addrspace(5)* [[TMP0]]) #[[ATTR5:[0-9]+]] -// CHECK-NEXT: [[TMP1:%.*]] = bitcast [4 x i8] addrspace(5)* [[S]] to i8 addrspace(5)* -// CHECK-NEXT: call void @llvm.memcpy.p5i8.p4i8.i64(i8 addrspace(5)* align 1 [[TMP1]], i8 addrspace(4)* align 1 getelementptr inbounds ([4 x i8], [4 x i8] addrspace(4)* @__const.test_printf_str_int.s, i32 0, i32 0), i64 4, i1 false) -// CHECK-NEXT: [[ARRAYDECAY:%.*]] = getelementptr inbounds [4 x i8], [4 x i8] addrspace(5)* [[S]], i64 0, i64 0 -// CHECK-NEXT: [[TMP2:%.*]] = load i32, i32 addrspace(5)* [[I_ADDR]], align 4, !tbaa [[TBAA8]] -// CHECK-NEXT: [[CALL:%.*]] = call i32 (i8 addrspace(4)*, ...) @printf(i8 addrspace(4)* noundef getelementptr inbounds ([6 x i8], [6 x i8] addrspace(4)* @.str.2, i64 0, i64 0), i8 addrspace(5)* noundef [[ARRAYDECAY]], i32 noundef [[TMP2]]) #[[ATTR4]] -// CHECK-NEXT: [[TMP3:%.*]] = bitcast [4 x i8] addrspace(5)* [[S]] to i8 addrspace(5)* -// CHECK-NEXT: call void @llvm.lifetime.end.p5i8(i64 4, i8 addrspace(5)* [[TMP3]]) #[[ATTR5]] +// CHECK-NEXT: store i32 [[I:%.*]], ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA8]] +// CHECK-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[S]]) #[[ATTR5:[0-9]+]] +// CHECK-NEXT: call void @llvm.memcpy.p5.p4.i64(ptr addrspace(5) align 1 [[S]], ptr addrspace(4) align 1 @__const.test_printf_str_int.s, i64 4, i1 false) +// CHECK-NEXT: [[ARRAYDECAY:%.*]] = getelementptr inbounds [4 x i8], ptr addrspace(5) [[S]], i64 0, i64 0 +// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA8]] +// CHECK-NEXT: [[CALL:%.*]] = call i32 (ptr addrspace(4), ...) @printf(ptr addrspace(4) noundef @.str.2, ptr addrspace(5) noundef [[ARRAYDECAY]], i32 noundef [[TMP2]]) #[[ATTR4]] +// CHECK-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[S]]) #[[ATTR5]] // CHECK-NEXT: ret void // __kernel void test_printf_str_int(int i) { diff --git a/clang/test/CodeGenOpenCL/as_type.cl b/clang/test/CodeGenOpenCL/as_type.cl index 355f026..afc76b7 100644 --- a/clang/test/CodeGenOpenCL/as_type.cl +++ b/clang/test/CodeGenOpenCL/as_type.cl @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -no-opaque-pointers %s -emit-llvm -triple spir-unknown-unknown -o - | FileCheck %s +// RUN: %clang_cc1 %s -emit-llvm -triple spir-unknown-unknown -o - | FileCheck %s typedef __attribute__(( ext_vector_type(3) )) char char3; typedef __attribute__(( ext_vector_type(4) )) char char4; @@ -67,29 +67,29 @@ int3 f8(char16 x) { return __builtin_astype(x, int3); } -//CHECK: define{{.*}} spir_func i32 addrspace(1)* @addr_cast(i32* noundef readnone %[[x:.*]]) -//CHECK: %[[cast:.*]] ={{.*}} addrspacecast i32* %[[x]] to i32 addrspace(1)* -//CHECK: ret i32 addrspace(1)* %[[cast]] +//CHECK: define{{.*}} spir_func ptr addrspace(1) @addr_cast(ptr noundef readnone %[[x:.*]]) +//CHECK: %[[cast:.*]] ={{.*}} addrspacecast ptr %[[x]] to ptr addrspace(1) +//CHECK: ret ptr addrspace(1) %[[cast]] global int* addr_cast(int *x) { return __builtin_astype(x, global int*); } -//CHECK: define{{.*}} spir_func i32 addrspace(1)* @int_to_ptr(i32 noundef %[[x:.*]]) -//CHECK: %[[cast:.*]] = inttoptr i32 %[[x]] to i32 addrspace(1)* -//CHECK: ret i32 addrspace(1)* %[[cast]] +//CHECK: define{{.*}} spir_func ptr addrspace(1) @int_to_ptr(i32 noundef %[[x:.*]]) +//CHECK: %[[cast:.*]] = inttoptr i32 %[[x]] to ptr addrspace(1) +//CHECK: ret ptr addrspace(1) %[[cast]] global int* int_to_ptr(int x) { return __builtin_astype(x, global int*); } -//CHECK: define{{.*}} spir_func i32 @ptr_to_int(i32* noundef %[[x:.*]]) -//CHECK: %[[cast:.*]] = ptrtoint i32* %[[x]] to i32 +//CHECK: define{{.*}} spir_func i32 @ptr_to_int(ptr noundef %[[x:.*]]) +//CHECK: %[[cast:.*]] = ptrtoint ptr %[[x]] to i32 //CHECK: ret i32 %[[cast]] int ptr_to_int(int *x) { return __builtin_astype(x, int); } -//CHECK: define{{.*}} spir_func <3 x i8> @ptr_to_char3(i32* noundef %[[x:.*]]) -//CHECK: %[[cast1:.*]] = ptrtoint i32* %[[x]] to i32 +//CHECK: define{{.*}} spir_func <3 x i8> @ptr_to_char3(ptr noundef %[[x:.*]]) +//CHECK: %[[cast1:.*]] = ptrtoint ptr %[[x]] to i32 //CHECK: %[[cast2:.*]] = bitcast i32 %[[cast1]] to <4 x i8> //CHECK: %[[astype:.*]] = shufflevector <4 x i8> %[[cast2]], <4 x i8> poison, <3 x i32> <i32 0, i32 1, i32 2> //CHECK: ret <3 x i8> %[[astype]] @@ -97,11 +97,11 @@ char3 ptr_to_char3(int *x) { return __builtin_astype(x, char3); } -//CHECK: define{{.*}} spir_func i32* @char3_to_ptr(<3 x i8> noundef %[[x:.*]]) +//CHECK: define{{.*}} spir_func ptr @char3_to_ptr(<3 x i8> noundef %[[x:.*]]) //CHECK: %[[astype:.*]] = shufflevector <3 x i8> %[[x]], <3 x i8> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 undef> //CHECK: %[[cast1:.*]] = bitcast <4 x i8> %[[astype]] to i32 -//CHECK: %[[cast2:.*]] = inttoptr i32 %[[cast1]] to i32* -//CHECK: ret i32* %[[cast2]] +//CHECK: %[[cast2:.*]] = inttoptr i32 %[[cast1]] to ptr +//CHECK: ret ptr %[[cast2]] int* char3_to_ptr(char3 x) { return __builtin_astype(x, int*); } diff --git a/clang/test/CodeGenOpenCL/atomics-cas-remarks-gfx90a.cl b/clang/test/CodeGenOpenCL/atomics-cas-remarks-gfx90a.cl index c62d58f..a5321ea 100644 --- a/clang/test/CodeGenOpenCL/atomics-cas-remarks-gfx90a.cl +++ b/clang/test/CodeGenOpenCL/atomics-cas-remarks-gfx90a.cl @@ -1,8 +1,8 @@ -// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL2.0 -O0 -triple=amdgcn-amd-amdhsa -target-cpu gfx90a \ +// RUN: %clang_cc1 %s -cl-std=CL2.0 -O0 -triple=amdgcn-amd-amdhsa -target-cpu gfx90a \ // RUN: -Rpass=atomic-expand -S -o - 2>&1 | \ // RUN: FileCheck %s --check-prefix=REMARK -// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL2.0 -O0 -triple=amdgcn-amd-amdhsa -target-cpu gfx90a \ +// RUN: %clang_cc1 %s -cl-std=CL2.0 -O0 -triple=amdgcn-amd-amdhsa -target-cpu gfx90a \ // RUN: -Rpass=atomic-expand -S -emit-llvm -o - 2>&1 | \ // RUN: FileCheck %s --check-prefix=GFX90A-CAS @@ -31,10 +31,10 @@ typedef enum memory_scope { // REMARK: remark: A compare and swap loop was generated for an atomic fadd operation at one-as memory scope [-Rpass=atomic-expand] // REMARK: remark: A compare and swap loop was generated for an atomic fadd operation at wavefront-one-as memory scope [-Rpass=atomic-expand] // GFX90A-CAS-LABEL: @atomic_cas -// GFX90A-CAS: atomicrmw fadd float addrspace(1)* {{.*}} syncscope("workgroup-one-as") monotonic -// GFX90A-CAS: atomicrmw fadd float addrspace(1)* {{.*}} syncscope("agent-one-as") monotonic -// GFX90A-CAS: atomicrmw fadd float addrspace(1)* {{.*}} syncscope("one-as") monotonic -// GFX90A-CAS: atomicrmw fadd float addrspace(1)* {{.*}} syncscope("wavefront-one-as") monotonic +// GFX90A-CAS: atomicrmw fadd ptr addrspace(1) {{.*}} syncscope("workgroup-one-as") monotonic +// GFX90A-CAS: atomicrmw fadd ptr addrspace(1) {{.*}} syncscope("agent-one-as") monotonic +// GFX90A-CAS: atomicrmw fadd ptr addrspace(1) {{.*}} syncscope("one-as") monotonic +// GFX90A-CAS: atomicrmw fadd ptr addrspace(1) {{.*}} syncscope("wavefront-one-as") monotonic float atomic_cas(__global atomic_float *d, float a) { float ret1 = __opencl_atomic_fetch_add(d, a, memory_order_relaxed, memory_scope_work_group); float ret2 = __opencl_atomic_fetch_add(d, a, memory_order_relaxed, memory_scope_device); diff --git a/clang/test/CodeGenOpenCL/atomics-unsafe-hw-remarks-gfx90a.cl b/clang/test/CodeGenOpenCL/atomics-unsafe-hw-remarks-gfx90a.cl index 61a24ac..1243745 100644 --- a/clang/test/CodeGenOpenCL/atomics-unsafe-hw-remarks-gfx90a.cl +++ b/clang/test/CodeGenOpenCL/atomics-unsafe-hw-remarks-gfx90a.cl @@ -1,8 +1,8 @@ -// RUN: %clang_cc1 -no-opaque-pointers -cl-std=CL2.0 -O0 -triple=amdgcn-amd-amdhsa -target-cpu gfx90a \ +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple=amdgcn-amd-amdhsa -target-cpu gfx90a \ // RUN: -Rpass=si-lower -munsafe-fp-atomics %s -S -emit-llvm -o - 2>&1 | \ // RUN: FileCheck %s --check-prefix=GFX90A-HW -// RUN: %clang_cc1 -no-opaque-pointers -cl-std=CL2.0 -O0 -triple=amdgcn-amd-amdhsa -target-cpu gfx90a \ +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple=amdgcn-amd-amdhsa -target-cpu gfx90a \ // RUN: -Rpass=si-lower -munsafe-fp-atomics %s -S -o - 2>&1 | \ // RUN: FileCheck %s --check-prefix=GFX90A-HW-REMARK @@ -34,9 +34,9 @@ typedef enum memory_scope { // GFX90A-HW-REMARK: global_atomic_add_f32 v0, v[0:1], v2, off glc // GFX90A-HW-REMARK: global_atomic_add_f32 v0, v[0:1], v2, off glc // GFX90A-HW-LABEL: @atomic_unsafe_hw -// GFX90A-HW: atomicrmw fadd float addrspace(1)* %{{.*}}, float %{{.*}} syncscope("workgroup-one-as") monotonic, align 4 -// GFX90A-HW: atomicrmw fadd float addrspace(1)* %{{.*}}, float %{{.*}} syncscope("agent-one-as") monotonic, align 4 -// GFX90A-HW: atomicrmw fadd float addrspace(1)* %{{.*}}, float %{{.*}} syncscope("wavefront-one-as") monotonic, align 4 +// GFX90A-HW: atomicrmw fadd ptr addrspace(1) %{{.*}}, float %{{.*}} syncscope("workgroup-one-as") monotonic, align 4 +// GFX90A-HW: atomicrmw fadd ptr addrspace(1) %{{.*}}, float %{{.*}} syncscope("agent-one-as") monotonic, align 4 +// GFX90A-HW: atomicrmw fadd ptr addrspace(1) %{{.*}}, float %{{.*}} syncscope("wavefront-one-as") monotonic, align 4 void atomic_unsafe_hw(__global atomic_float *d, float a) { float ret1 = __opencl_atomic_fetch_add(d, a, memory_order_relaxed, memory_scope_work_group); float ret2 = __opencl_atomic_fetch_add(d, a, memory_order_relaxed, memory_scope_device); diff --git a/clang/test/CodeGenOpenCL/bool_cast.cl b/clang/test/CodeGenOpenCL/bool_cast.cl index 9db3d4a2..6e6538f 100644 --- a/clang/test/CodeGenOpenCL/bool_cast.cl +++ b/clang/test/CodeGenOpenCL/bool_cast.cl @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -no-opaque-pointers %s -triple x86_64-unknown-linux-gnu -emit-llvm -o - -O0 | FileCheck %s +// RUN: %clang_cc1 %s -triple x86_64-unknown-linux-gnu -emit-llvm -o - -O0 | FileCheck %s typedef unsigned char uchar4 __attribute((ext_vector_type(4))); typedef unsigned int int4 __attribute((ext_vector_type(4))); @@ -8,24 +8,24 @@ typedef float float4 __attribute((ext_vector_type(4))); void kernel ker() { bool t = true; int4 vec4 = (int4)t; -// CHECK: {{%.*}} = load i8, i8* %t, align 1 +// CHECK: {{%.*}} = load i8, ptr %t, align 1 // CHECK: {{%.*}} = trunc i8 {{%.*}} to i1 // CHECK: {{%.*}} = sext i1 {{%.*}} to i32 // CHECK: {{%.*}} = insertelement <4 x i32> poison, i32 {{%.*}}, i32 0 // CHECK: {{%.*}} = shufflevector <4 x i32> {{%.*}}, <4 x i32> poison, <4 x i32> zeroinitializer -// CHECK: store <4 x i32> {{%.*}}, <4 x i32>* %vec4, align 16 +// CHECK: store <4 x i32> {{%.*}}, ptr %vec4, align 16 int i = (int)t; -// CHECK: {{%.*}} = load i8, i8* %t, align 1 +// CHECK: {{%.*}} = load i8, ptr %t, align 1 // CHECK: {{%.*}} = trunc i8 {{%.*}} to i1 // CHECK: {{%.*}} = zext i1 {{%.*}} to i32 -// CHECK: store i32 {{%.*}}, i32* %i, align 4 +// CHECK: store i32 {{%.*}}, ptr %i, align 4 uchar4 vc; vc = (uchar4)true; -// CHECK: store <4 x i8> <i8 -1, i8 -1, i8 -1, i8 -1>, <4 x i8>* %vc, align 4 +// CHECK: store <4 x i8> <i8 -1, i8 -1, i8 -1, i8 -1>, ptr %vc, align 4 unsigned char c; c = (unsigned char)true; -// CHECK: store i8 1, i8* %c, align 1 +// CHECK: store i8 1, ptr %c, align 1 float4 vf; vf = (float4)true; diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl index 0a752a4..da989ec 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl @@ -1,8 +1,8 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -no-opaque-pointers -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu hawaii -S -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -no-opaque-pointers -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu fiji -S -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -no-opaque-pointers -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -no-opaque-pointers -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu hawaii -S -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu fiji -S -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s typedef unsigned int uint; typedef unsigned long ulong; @@ -36,29 +36,27 @@ void test_s_memtime(global ulong* out) } // CHECK-LABEL: @test_is_shared( -// CHECK: [[CAST:%[0-9]+]] = bitcast i32* %{{[0-9]+}} to i8* -// CHECK: call i1 @llvm.amdgcn.is.shared(i8* [[CAST]] +// CHECK: call i1 @llvm.amdgcn.is.shared(ptr %{{[0-9]+}} int test_is_shared(const int* ptr) { return __builtin_amdgcn_is_shared(ptr); } // CHECK-LABEL: @test_is_private( -// CHECK: [[CAST:%[0-9]+]] = bitcast i32* %{{[0-9]+}} to i8* -// CHECK: call i1 @llvm.amdgcn.is.private(i8* [[CAST]] +// CHECK: call i1 @llvm.amdgcn.is.private(ptr %{{[0-9]+}} int test_is_private(const int* ptr) { return __builtin_amdgcn_is_private(ptr); } // CHECK-LABEL: @test_is_shared_global( -// CHECK: [[CAST:%[0-9]+]] = addrspacecast i32 addrspace(1)* %{{[0-9]+}} to i8* -// CHECK: call i1 @llvm.amdgcn.is.shared(i8* [[CAST]] +// CHECK: [[CAST:%[0-9]+]] = addrspacecast ptr addrspace(1) %{{[0-9]+}} to ptr +// CHECK: call i1 @llvm.amdgcn.is.shared(ptr [[CAST]] int test_is_shared_global(const global int* ptr) { return __builtin_amdgcn_is_shared(ptr); } // CHECK-LABEL: @test_is_private_global( -// CHECK: [[CAST:%[0-9]+]] = addrspacecast i32 addrspace(1)* %{{[0-9]+}} to i8* -// CHECK: call i1 @llvm.amdgcn.is.private(i8* [[CAST]] +// CHECK: [[CAST:%[0-9]+]] = addrspacecast ptr addrspace(1) %{{[0-9]+}} to ptr +// CHECK: call i1 @llvm.amdgcn.is.private(ptr [[CAST]] int test_is_private_global(const global int* ptr) { return __builtin_amdgcn_is_private(ptr); } diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl index c7437d7..590670b 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl @@ -1,8 +1,8 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-unknown-unknown -target-cpu tonga -S -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-unknown-unknown -target-cpu gfx900 -S -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-unknown-unknown -target-cpu gfx1012 -S -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu tonga -S -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -S -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1012 -S -emit-llvm -o - %s | FileCheck %s #pragma OPENCL EXTENSION cl_khr_fp16 : enable @@ -115,19 +115,19 @@ void test_update_dpp(global int* out, int arg1, int arg2) } // CHECK-LABEL: @test_ds_fadd -// CHECK: call float @llvm.amdgcn.ds.fadd.f32(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false) +// CHECK: call float @llvm.amdgcn.ds.fadd.f32(ptr addrspace(3) %out, float %src, i32 0, i32 0, i1 false) void test_ds_faddf(local float *out, float src) { *out = __builtin_amdgcn_ds_faddf(out, src, 0, 0, false); } // CHECK-LABEL: @test_ds_fmin -// CHECK: call float @llvm.amdgcn.ds.fmin.f32(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false) +// CHECK: call float @llvm.amdgcn.ds.fmin.f32(ptr addrspace(3) %out, float %src, i32 0, i32 0, i1 false) void test_ds_fminf(local float *out, float src) { *out = __builtin_amdgcn_ds_fminf(out, src, 0, 0, false); } // CHECK-LABEL: @test_ds_fmax -// CHECK: call float @llvm.amdgcn.ds.fmax.f32(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false) +// CHECK: call float @llvm.amdgcn.ds.fmax.f32(ptr addrspace(3) %out, float %src, i32 0, i32 0, i1 false) void test_ds_fmaxf(local float *out, float src) { *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 0, false); } diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index 9a20dba..82cd317 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -1,5 +1,5 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -no-opaque-pointers -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu tahiti -S -emit-llvm -o - %s | FileCheck -enable-var-scope %s +// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu tahiti -S -emit-llvm -o - %s | FileCheck -enable-var-scope %s #pragma OPENCL EXTENSION cl_khr_fp64 : enable @@ -516,28 +516,28 @@ void test_read_exec_hi(global uint* out) { } // CHECK-LABEL: @test_dispatch_ptr -// CHECK: call align 4 dereferenceable(64) i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr() +// CHECK: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() void test_dispatch_ptr(__constant unsigned char ** out) { *out = __builtin_amdgcn_dispatch_ptr(); } // CHECK-LABEL: @test_queue_ptr -// CHECK: call i8 addrspace(4)* @llvm.amdgcn.queue.ptr() +// CHECK: call ptr addrspace(4) @llvm.amdgcn.queue.ptr() void test_queue_ptr(__constant unsigned char ** out) { *out = __builtin_amdgcn_queue_ptr(); } // CHECK-LABEL: @test_kernarg_segment_ptr -// CHECK: call i8 addrspace(4)* @llvm.amdgcn.kernarg.segment.ptr() +// CHECK: call ptr addrspace(4) @llvm.amdgcn.kernarg.segment.ptr() void test_kernarg_segment_ptr(__constant unsigned char ** out) { *out = __builtin_amdgcn_kernarg_segment_ptr(); } // CHECK-LABEL: @test_implicitarg_ptr -// CHECK: call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr() +// CHECK: call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() void test_implicitarg_ptr(__constant unsigned char ** out) { *out = __builtin_amdgcn_implicitarg_ptr(); @@ -583,13 +583,13 @@ void test_get_local_id(int d, global int *out) } // CHECK-LABEL: @test_get_workgroup_size( -// CHECK: call align 4 dereferenceable(64) i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr() -// CHECK: getelementptr i8, i8 addrspace(4)* %{{.*}}, i64 4 -// CHECK: load i16, i16 addrspace(4)* %{{.*}}, align 4, !range [[$WS_RANGE:![0-9]*]], !invariant.load -// CHECK: getelementptr i8, i8 addrspace(4)* %{{.*}}, i64 6 -// CHECK: load i16, i16 addrspace(4)* %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load -// CHECK: getelementptr i8, i8 addrspace(4)* %{{.*}}, i64 8 -// CHECK: load i16, i16 addrspace(4)* %{{.*}}, align 4, !range [[$WS_RANGE:![0-9]*]], !invariant.load +// CHECK: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() +// CHECK: getelementptr i8, ptr addrspace(4) %{{.*}}, i64 4 +// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 4, !range [[$WS_RANGE:![0-9]*]], !invariant.load +// CHECK: getelementptr i8, ptr addrspace(4) %{{.*}}, i64 6 +// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load +// CHECK: getelementptr i8, ptr addrspace(4) %{{.*}}, i64 8 +// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 4, !range [[$WS_RANGE:![0-9]*]], !invariant.load void test_get_workgroup_size(int d, global int *out) { switch (d) { @@ -601,13 +601,13 @@ void test_get_workgroup_size(int d, global int *out) } // CHECK-LABEL: @test_get_grid_size( -// CHECK: call align 4 dereferenceable(64) i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr() -// CHECK: getelementptr i8, i8 addrspace(4)* %{{.*}}, i64 12 -// CHECK: load i32, i32 addrspace(4)* %{{.*}}, align 4, !invariant.load -// CHECK: getelementptr i8, i8 addrspace(4)* %{{.*}}, i64 16 -// CHECK: load i32, i32 addrspace(4)* %{{.*}}, align 4, !invariant.load -// CHECK: getelementptr i8, i8 addrspace(4)* %{{.*}}, i64 20 -// CHECK: load i32, i32 addrspace(4)* %{{.*}}, align 4, !invariant.load +// CHECK: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() +// CHECK: getelementptr i8, ptr addrspace(4) %{{.*}}, i64 12 +// CHECK: load i32, ptr addrspace(4) %{{.*}}, align 4, !invariant.load +// CHECK: getelementptr i8, ptr addrspace(4) %{{.*}}, i64 16 +// CHECK: load i32, ptr addrspace(4) %{{.*}}, align 4, !invariant.load +// CHECK: getelementptr i8, ptr addrspace(4) %{{.*}}, i64 20 +// CHECK: load i32, ptr addrspace(4) %{{.*}}, align 4, !invariant.load void test_get_grid_size(int d, global int *out) { switch (d) { @@ -633,13 +633,13 @@ void test_s_getpc(global ulong* out) } // CHECK-LABEL: @test_ds_append_lds( -// CHECK: call i32 @llvm.amdgcn.ds.append.p3i32(i32 addrspace(3)* %ptr, i1 false) +// CHECK: call i32 @llvm.amdgcn.ds.append.p3(ptr addrspace(3) %ptr, i1 false) kernel void test_ds_append_lds(global int* out, local int* ptr) { *out = __builtin_amdgcn_ds_append(ptr); } // CHECK-LABEL: @test_ds_consume_lds( -// CHECK: call i32 @llvm.amdgcn.ds.consume.p3i32(i32 addrspace(3)* %ptr, i1 false) +// CHECK: call i32 @llvm.amdgcn.ds.consume.p3(ptr addrspace(3) %ptr, i1 false) kernel void test_ds_consume_lds(global int* out, local int* ptr) { *out = __builtin_amdgcn_ds_consume(ptr); } diff --git a/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx8.cl b/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx8.cl index e4d85f7..f9782c1 100644 --- a/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx8.cl +++ b/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx8.cl @@ -1,12 +1,12 @@ -// RUN: %clang_cc1 -no-opaque-pointers -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx810 \ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx810 \ // RUN: %s -S -emit-llvm -o - | FileCheck %s -// RUN: %clang_cc1 -no-opaque-pointers -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx810 \ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx810 \ // RUN: -S -o - %s | FileCheck -check-prefix=GFX8 %s // REQUIRES: amdgpu-registered-target // CHECK-LABEL: test_fadd_local -// CHECK: call float @llvm.amdgcn.ds.fadd.f32(float addrspace(3)* %{{.*}}, float %{{.*}}, i32 0, i32 0, i1 false) +// CHECK: call float @llvm.amdgcn.ds.fadd.f32(ptr addrspace(3) %{{.*}}, float %{{.*}}, i32 0, i32 0, i1 false) // GFX8-LABEL: test_fadd_local$local: // GFX8: ds_add_rtn_f32 v2, v0, v1 // GFX8: s_endpgm diff --git a/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx90a.cl b/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx90a.cl index f078f4e..c582c7f 100644 --- a/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx90a.cl +++ b/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx90a.cl @@ -1,7 +1,7 @@ -// RUN: %clang_cc1 -no-opaque-pointers -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \ // RUN: %s -S -emit-llvm -o - | FileCheck %s -check-prefix=CHECK -// RUN: %clang_cc1 -no-opaque-pointers -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \ // RUN: -S -o - %s | FileCheck -check-prefix=GFX90A %s // REQUIRES: amdgpu-registered-target @@ -9,7 +9,7 @@ typedef half __attribute__((ext_vector_type(2))) half2; // CHECK-LABEL: test_global_add_f64 -// CHECK: call double @llvm.amdgcn.global.atomic.fadd.f64.p1f64.f64(double addrspace(1)* %{{.*}}, double %{{.*}}) +// CHECK: call double @llvm.amdgcn.global.atomic.fadd.f64.p1.f64(ptr addrspace(1) %{{.*}}, double %{{.*}}) // GFX90A-LABEL: test_global_add_f64$local: // GFX90A: global_atomic_add_f64 void test_global_add_f64(__global double *addr, double x) { @@ -18,7 +18,7 @@ void test_global_add_f64(__global double *addr, double x) { } // CHECK-LABEL: test_global_add_half2 -// CHECK: call <2 x half> @llvm.amdgcn.global.atomic.fadd.v2f16.p1v2f16.v2f16(<2 x half> addrspace(1)* %{{.*}}, <2 x half> %{{.*}}) +// CHECK: call <2 x half> @llvm.amdgcn.global.atomic.fadd.v2f16.p1.v2f16(ptr addrspace(1) %{{.*}}, <2 x half> %{{.*}}) // GFX90A-LABEL: test_global_add_half2 // GFX90A: global_atomic_pk_add_f16 v2, v[0:1], v2, off glc void test_global_add_half2(__global half2 *addr, half2 x) { @@ -27,7 +27,7 @@ void test_global_add_half2(__global half2 *addr, half2 x) { } // CHECK-LABEL: test_global_global_min_f64 -// CHECK: call double @llvm.amdgcn.global.atomic.fmin.f64.p1f64.f64(double addrspace(1)* %{{.*}}, double %{{.*}}) +// CHECK: call double @llvm.amdgcn.global.atomic.fmin.f64.p1.f64(ptr addrspace(1) %{{.*}}, double %{{.*}}) // GFX90A-LABEL: test_global_global_min_f64$local // GFX90A: global_atomic_min_f64 void test_global_global_min_f64(__global double *addr, double x){ @@ -36,7 +36,7 @@ void test_global_global_min_f64(__global double *addr, double x){ } // CHECK-LABEL: test_global_max_f64 -// CHECK: call double @llvm.amdgcn.global.atomic.fmax.f64.p1f64.f64(double addrspace(1)* %{{.*}}, double %{{.*}}) +// CHECK: call double @llvm.amdgcn.global.atomic.fmax.f64.p1.f64(ptr addrspace(1) %{{.*}}, double %{{.*}}) // GFX90A-LABEL: test_global_max_f64$local // GFX90A: global_atomic_max_f64 void test_global_max_f64(__global double *addr, double x){ @@ -45,7 +45,7 @@ void test_global_max_f64(__global double *addr, double x){ } // CHECK-LABEL: test_flat_add_local_f64 -// CHECK: call double @llvm.amdgcn.flat.atomic.fadd.f64.p3f64.f64(double addrspace(3)* %{{.*}}, double %{{.*}}) +// CHECK: call double @llvm.amdgcn.flat.atomic.fadd.f64.p3.f64(ptr addrspace(3) %{{.*}}, double %{{.*}}) // GFX90A-LABEL: test_flat_add_local_f64$local // GFX90A: ds_add_rtn_f64 void test_flat_add_local_f64(__local double *addr, double x){ @@ -54,7 +54,7 @@ void test_flat_add_local_f64(__local double *addr, double x){ } // CHECK-LABEL: test_flat_global_add_f64 -// CHECK: call double @llvm.amdgcn.flat.atomic.fadd.f64.p1f64.f64(double addrspace(1)* %{{.*}}, double %{{.*}}) +// CHECK: call double @llvm.amdgcn.flat.atomic.fadd.f64.p1.f64(ptr addrspace(1) %{{.*}}, double %{{.*}}) // GFX90A-LABEL: test_flat_global_add_f64$local // GFX90A: global_atomic_add_f64 void test_flat_global_add_f64(__global double *addr, double x){ @@ -63,7 +63,7 @@ void test_flat_global_add_f64(__global double *addr, double x){ } // CHECK-LABEL: test_flat_min_flat_f64 -// CHECK: call double @llvm.amdgcn.flat.atomic.fmin.f64.p0f64.f64(double* %{{.*}}, double %{{.*}}) +// CHECK: call double @llvm.amdgcn.flat.atomic.fmin.f64.p0.f64(ptr %{{.*}}, double %{{.*}}) // GFX90A-LABEL: test_flat_min_flat_f64$local // GFX90A: flat_atomic_min_f64 void test_flat_min_flat_f64(__generic double *addr, double x){ @@ -72,7 +72,7 @@ void test_flat_min_flat_f64(__generic double *addr, double x){ } // CHECK-LABEL: test_flat_global_min_f64 -// CHECK: call double @llvm.amdgcn.flat.atomic.fmin.f64.p1f64.f64(double addrspace(1)* %{{.*}}, double %{{.*}}) +// CHECK: call double @llvm.amdgcn.flat.atomic.fmin.f64.p1.f64(ptr addrspace(1) %{{.*}}, double %{{.*}}) // GFX90A: test_flat_global_min_f64$local // GFX90A: global_atomic_min_f64 void test_flat_global_min_f64(__global double *addr, double x){ @@ -81,7 +81,7 @@ void test_flat_global_min_f64(__global double *addr, double x){ } // CHECK-LABEL: test_flat_max_flat_f64 -// CHECK: call double @llvm.amdgcn.flat.atomic.fmax.f64.p0f64.f64(double* %{{.*}}, double %{{.*}}) +// CHECK: call double @llvm.amdgcn.flat.atomic.fmax.f64.p0.f64(ptr %{{.*}}, double %{{.*}}) // GFX90A-LABEL: test_flat_max_flat_f64$local // GFX90A: flat_atomic_max_f64 void test_flat_max_flat_f64(__generic double *addr, double x){ @@ -90,7 +90,7 @@ void test_flat_max_flat_f64(__generic double *addr, double x){ } // CHECK-LABEL: test_flat_global_max_f64 -// CHECK: call double @llvm.amdgcn.flat.atomic.fmax.f64.p1f64.f64(double addrspace(1)* %{{.*}}, double %{{.*}}) +// CHECK: call double @llvm.amdgcn.flat.atomic.fmax.f64.p1.f64(ptr addrspace(1) %{{.*}}, double %{{.*}}) // GFX90A-LABEL: test_flat_global_max_f64$local // GFX90A: global_atomic_max_f64 void test_flat_global_max_f64(__global double *addr, double x){ @@ -99,7 +99,7 @@ void test_flat_global_max_f64(__global double *addr, double x){ } // CHECK-LABEL: test_ds_add_local_f64 -// CHECK: call double @llvm.amdgcn.ds.fadd.f64(double addrspace(3)* %{{.*}}, double %{{.*}}, +// CHECK: call double @llvm.amdgcn.ds.fadd.f64(ptr addrspace(3) %{{.*}}, double %{{.*}}, // GFX90A: test_ds_add_local_f64$local // GFX90A: ds_add_rtn_f64 void test_ds_add_local_f64(__local double *addr, double x){ @@ -108,7 +108,7 @@ void test_ds_add_local_f64(__local double *addr, double x){ } // CHECK-LABEL: test_ds_addf_local_f32 -// CHECK: call float @llvm.amdgcn.ds.fadd.f32(float addrspace(3)* %{{.*}}, float %{{.*}}, +// CHECK: call float @llvm.amdgcn.ds.fadd.f32(ptr addrspace(3) %{{.*}}, float %{{.*}}, // GFX90A-LABEL: test_ds_addf_local_f32$local // GFX90A: ds_add_rtn_f32 void test_ds_addf_local_f32(__local float *addr, float x){ diff --git a/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx940.cl b/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx940.cl index 8f22f82..12b8593 100644 --- a/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx940.cl +++ b/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx940.cl @@ -1,7 +1,7 @@ -// RUN: %clang_cc1 -no-opaque-pointers -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx940 \ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx940 \ // RUN: %s -S -emit-llvm -o - | FileCheck %s -// RUN: %clang_cc1 -no-opaque-pointers -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx940 \ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx940 \ // RUN: -S -o - %s | FileCheck -check-prefix=GFX940 %s // REQUIRES: amdgpu-registered-target @@ -10,7 +10,7 @@ typedef half __attribute__((ext_vector_type(2))) half2; typedef short __attribute__((ext_vector_type(2))) short2; // CHECK-LABEL: test_flat_add_f32 -// CHECK: call float @llvm.amdgcn.flat.atomic.fadd.f32.p0f32.f32(float* %{{.*}}, float %{{.*}}) +// CHECK: call float @llvm.amdgcn.flat.atomic.fadd.f32.p0.f32(ptr %{{.*}}, float %{{.*}}) // GFX940-LABEL: test_flat_add_f32 // GFX940: flat_atomic_add_f32 half2 test_flat_add_f32(__generic float *addr, float x) { @@ -18,7 +18,7 @@ half2 test_flat_add_f32(__generic float *addr, float x) { } // CHECK-LABEL: test_flat_add_2f16 -// CHECK: call <2 x half> @llvm.amdgcn.flat.atomic.fadd.v2f16.p0v2f16.v2f16(<2 x half>* %{{.*}}, <2 x half> %{{.*}}) +// CHECK: call <2 x half> @llvm.amdgcn.flat.atomic.fadd.v2f16.p0.v2f16(ptr %{{.*}}, <2 x half> %{{.*}}) // GFX940-LABEL: test_flat_add_2f16 // GFX940: flat_atomic_pk_add_f16 half2 test_flat_add_2f16(__generic half2 *addr, half2 x) { @@ -26,7 +26,7 @@ half2 test_flat_add_2f16(__generic half2 *addr, half2 x) { } // CHECK-LABEL: test_flat_add_2bf16 -// CHECK: call <2 x i16> @llvm.amdgcn.flat.atomic.fadd.v2bf16.p0v2i16(<2 x i16>* %{{.*}}, <2 x i16> %{{.*}}) +// CHECK: call <2 x i16> @llvm.amdgcn.flat.atomic.fadd.v2bf16.p0(ptr %{{.*}}, <2 x i16> %{{.*}}) // GFX940-LABEL: test_flat_add_2bf16 // GFX940: flat_atomic_pk_add_bf16 short2 test_flat_add_2bf16(__generic short2 *addr, short2 x) { @@ -34,7 +34,7 @@ short2 test_flat_add_2bf16(__generic short2 *addr, short2 x) { } // CHECK-LABEL: test_global_add_2bf16 -// CHECK: call <2 x i16> @llvm.amdgcn.global.atomic.fadd.v2bf16.p1v2i16(<2 x i16> addrspace(1)* %{{.*}}, <2 x i16> %{{.*}}) +// CHECK: call <2 x i16> @llvm.amdgcn.global.atomic.fadd.v2bf16.p1(ptr addrspace(1) %{{.*}}, <2 x i16> %{{.*}}) // GFX940-LABEL: test_global_add_2bf16 // GFX940: global_atomic_pk_add_bf16 short2 test_global_add_2bf16(__global short2 *addr, short2 x) { @@ -42,7 +42,7 @@ short2 test_global_add_2bf16(__global short2 *addr, short2 x) { } // CHECK-LABEL: test_local_add_2bf16 -// CHECK: call <2 x i16> @llvm.amdgcn.ds.fadd.v2bf16(<2 x i16> addrspace(3)* %{{.*}}, <2 x i16> % +// CHECK: call <2 x i16> @llvm.amdgcn.ds.fadd.v2bf16(ptr addrspace(3) %{{.*}}, <2 x i16> % // GFX940-LABEL: test_local_add_2bf16 // GFX940: ds_pk_add_rtn_bf16 short2 test_local_add_2bf16(__local short2 *addr, short2 x) { diff --git a/clang/test/CodeGenOpenCL/builtins-generic-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-generic-amdgcn.cl index 7ead1f0..46c044a0 100644 --- a/clang/test/CodeGenOpenCL/builtins-generic-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-generic-amdgcn.cl @@ -1,5 +1,5 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -no-opaque-pointers -Wno-error=int-conversion -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -Wno-error=int-conversion -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s // CHECK-LABEL: @test_builtin_clz( // CHECK: tail call i32 @llvm.ctlz.i32(i32 %a, i1 true) @@ -15,7 +15,7 @@ void test_builtin_clzl(global long* out, long a) *out = __builtin_clzl(a); } -// CHECK: tail call i8 addrspace(5)* @llvm.frameaddress.p5i8(i32 0) +// CHECK: tail call ptr addrspace(5) @llvm.frameaddress.p5(i32 0) void test_builtin_frame_address(int *out) { *out = __builtin_frame_address(0); } diff --git a/clang/test/CodeGenOpenCL/builtins-r600.cl b/clang/test/CodeGenOpenCL/builtins-r600.cl index 274de15..7ceb8d6 100644 --- a/clang/test/CodeGenOpenCL/builtins-r600.cl +++ b/clang/test/CodeGenOpenCL/builtins-r600.cl @@ -1,5 +1,5 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -no-opaque-pointers -triple r600-unknown-unknown -target-cpu cypress -S -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple r600-unknown-unknown -target-cpu cypress -S -emit-llvm -o - %s | FileCheck %s // CHECK-LABEL: @test_recipsqrt_ieee_f32 // CHECK: call float @llvm.r600.recipsqrt.ieee.f32 @@ -18,7 +18,7 @@ void test_recipsqrt_ieee_f64(global double* out, double a) #endif // CHECK-LABEL: @test_implicitarg_ptr -// CHECK: call i8 addrspace(7)* @llvm.r600.implicitarg.ptr() +// CHECK: call ptr addrspace(7) @llvm.r600.implicitarg.ptr() void test_implicitarg_ptr(__attribute__((address_space(7))) unsigned char ** out) { *out = __builtin_r600_implicitarg_ptr(); diff --git a/clang/test/CodeGenOpenCL/byval.cl b/clang/test/CodeGenOpenCL/byval.cl index bd994dd..6e734d7 100644 --- a/clang/test/CodeGenOpenCL/byval.cl +++ b/clang/test/CodeGenOpenCL/byval.cl @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -no-opaque-pointers -emit-llvm -o - -triple amdgcn %s | FileCheck %s +// RUN: %clang_cc1 -emit-llvm -o - -triple amdgcn %s | FileCheck %s struct A { int x[100]; @@ -8,8 +8,8 @@ int f(struct A a); int g() { struct A a; - // CHECK: call i32 @f(%struct.A addrspace(5)* noundef byval{{.*}}%a) + // CHECK: call i32 @f(ptr addrspace(5) noundef byval{{.*}}%a) return f(a); } -// CHECK: declare i32 @f(%struct.A addrspace(5)* noundef byval{{.*}}) +// CHECK: declare i32 @f(ptr addrspace(5) noundef byval{{.*}}) diff --git a/clang/test/CodeGenOpenCL/constant-addr-space-globals.cl b/clang/test/CodeGenOpenCL/constant-addr-space-globals.cl index adb9252..c38d3ed 100644 --- a/clang/test/CodeGenOpenCL/constant-addr-space-globals.cl +++ b/clang/test/CodeGenOpenCL/constant-addr-space-globals.cl @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -no-opaque-pointers %s -triple "spir64-unknown-unknown" -cl-opt-disable -ffake-address-space-map -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple "spir64-unknown-unknown" -cl-opt-disable -ffake-address-space-map -emit-llvm -o - | FileCheck %s // CHECK: @array ={{.*}} addrspace({{[0-9]+}}) constant __constant float array[2] = {0.0f, 1.0f}; @@ -26,6 +26,6 @@ kernel void k(void) { constant int var1 = 1; - // CHECK: call spir_func void @foo(i32 addrspace(2)* noundef @k.var1, i32 addrspace(2)* noundef getelementptr inbounds ([3 x i32], [3 x i32] addrspace(2)* @k.arr1, i64 0, i64 0) + // CHECK: call spir_func void @foo(ptr addrspace(2) noundef @k.var1, ptr addrspace(2) noundef @k.arr1 foo(&var1, arr1, arr2, arr3); } diff --git a/clang/test/CodeGenOpenCL/event_t.cl b/clang/test/CodeGenOpenCL/event_t.cl index b94332b..bb10414 100644 --- a/clang/test/CodeGenOpenCL/event_t.cl +++ b/clang/test/CodeGenOpenCL/event_t.cl @@ -1,14 +1,14 @@ -// RUN: %clang_cc1 -no-opaque-pointers %s -emit-llvm -o - -O0 | FileCheck %s +// RUN: %clang_cc1 %s -emit-llvm -o - -O0 | FileCheck %s void foo(event_t evt); void kernel ker() { event_t e; -// CHECK: alloca %opencl.event_t*, +// CHECK: alloca ptr, foo(e); -// CHECK: call {{.*}}void @foo(%opencl.event_t* % +// CHECK: call {{.*}}void @foo(ptr % foo(0); -// CHECK: call {{.*}}void @foo(%opencl.event_t* null) +// CHECK: call {{.*}}void @foo(ptr null) foo((event_t)0); -// CHECK: call {{.*}}void @foo(%opencl.event_t* null) +// CHECK: call {{.*}}void @foo(ptr null) } diff --git a/clang/test/CodeGenOpenCL/images.cl b/clang/test/CodeGenOpenCL/images.cl index 91867ba..f15397a 100644 --- a/clang/test/CodeGenOpenCL/images.cl +++ b/clang/test/CodeGenOpenCL/images.cl @@ -1,12 +1,12 @@ -// RUN: %clang_cc1 -no-opaque-pointers %s -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -o - | FileCheck %s -// RUN: %clang_cc1 -no-opaque-pointers %s -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -o - -cl-std=clc++ | FileCheck %s +// RUN: %clang_cc1 %s -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -o - -cl-std=clc++ | FileCheck %s __attribute__((overloadable)) void read_image(read_only image1d_t img_ro); __attribute__((overloadable)) void read_image(write_only image1d_t img_wo); kernel void test_read_image(read_only image1d_t img_ro, write_only image1d_t img_wo) { - // CHECK: call void @_Z10read_image14ocl_image1d_ro(%opencl.image1d_ro_t* %{{[0-9]+}}) + // CHECK: call void @_Z10read_image14ocl_image1d_ro(ptr %{{[0-9]+}}) read_image(img_ro); - // CHECK: call void @_Z10read_image14ocl_image1d_wo(%opencl.image1d_wo_t* %{{[0-9]+}}) + // CHECK: call void @_Z10read_image14ocl_image1d_wo(ptr %{{[0-9]+}}) read_image(img_wo); } diff --git a/clang/test/CodeGenOpenCL/kernel-param-alignment.cl b/clang/test/CodeGenOpenCL/kernel-param-alignment.cl index d4e6628..c1a1740 100644 --- a/clang/test/CodeGenOpenCL/kernel-param-alignment.cl +++ b/clang/test/CodeGenOpenCL/kernel-param-alignment.cl @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL1.2 -emit-llvm -triple x86_64-unknown-unknown -o - | FileCheck %s +// RUN: %clang_cc1 %s -cl-std=CL1.2 -emit-llvm -triple x86_64-unknown-unknown -o - | FileCheck %s // Test that pointer arguments to kernels are assumed to be ABI aligned. @@ -17,10 +17,10 @@ kernel void test( global void *v, global struct packed *p) { // CHECK-LABEL: spir_kernel void @test( -// CHECK-SAME: i32* nocapture noundef align 4 %i32, -// CHECK-SAME: i64* nocapture noundef align 8 %i64, -// CHECK-SAME: <4 x i32>* nocapture noundef align 16 %v4i32, -// CHECK-SAME: <2 x float>* nocapture noundef align 8 %v2f32, -// CHECK-SAME: i8* nocapture noundef %v, -// CHECK-SAME: %struct.packed* nocapture noundef align 1 %p) +// CHECK-SAME: ptr nocapture noundef align 4 %i32, +// CHECK-SAME: ptr nocapture noundef align 8 %i64, +// CHECK-SAME: ptr nocapture noundef align 16 %v4i32, +// CHECK-SAME: ptr nocapture noundef align 8 %v2f32, +// CHECK-SAME: ptr nocapture noundef %v, +// CHECK-SAME: ptr nocapture noundef align 1 %p) } diff --git a/clang/test/CodeGenOpenCL/lifetime.cl b/clang/test/CodeGenOpenCL/lifetime.cl index c016835..45df0c4 100644 --- a/clang/test/CodeGenOpenCL/lifetime.cl +++ b/clang/test/CodeGenOpenCL/lifetime.cl @@ -1,5 +1,5 @@ -// RUN: %clang_cc1 -no-opaque-pointers -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -no-opaque-pointers -emit-llvm -o - -triple amdgcn %s | FileCheck %s -check-prefix=AMDGCN +// RUN: %clang_cc1 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -emit-llvm -o - -triple amdgcn %s | FileCheck %s -check-prefix=AMDGCN void use(char *a); @@ -9,7 +9,7 @@ __attribute__((always_inline)) void helper_no_markers() { } void lifetime_test() { -// CHECK: @llvm.lifetime.start.p0i -// AMDGCN: @llvm.lifetime.start.p5i +// CHECK: @llvm.lifetime.start.p0 +// AMDGCN: @llvm.lifetime.start.p5 helper_no_markers(); } diff --git a/clang/test/CodeGenOpenCL/memcpy.cl b/clang/test/CodeGenOpenCL/memcpy.cl index 5255a2f..06362ff 100644 --- a/clang/test/CodeGenOpenCL/memcpy.cl +++ b/clang/test/CodeGenOpenCL/memcpy.cl @@ -1,8 +1,8 @@ -// RUN: %clang_cc1 -no-opaque-pointers %s -ffake-address-space-map -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck %s // CHECK-LABEL: @test // CHECK-NOT: addrspacecast -// CHECK: call void @llvm.memcpy.p1i8.p2i8 +// CHECK: call void @llvm.memcpy.p1.p2 kernel void test(global float *g, constant float *c) { __builtin_memcpy(g, c, 32); } diff --git a/clang/test/CodeGenOpenCL/null_queue.cl b/clang/test/CodeGenOpenCL/null_queue.cl index aaa9457..d6b3123 100644 --- a/clang/test/CodeGenOpenCL/null_queue.cl +++ b/clang/test/CodeGenOpenCL/null_queue.cl @@ -1,11 +1,11 @@ -// RUN: %clang_cc1 -no-opaque-pointers -O0 -cl-std=CL2.0 -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -emit-llvm %s -o - | FileCheck %s extern queue_t get_default_queue(void); bool compare(void) { return 0 == get_default_queue() && get_default_queue() == 0; - // CHECK: icmp eq %opencl.queue_t* null, %{{.*}} - // CHECK: icmp eq %opencl.queue_t* %{{.*}}, null + // CHECK: icmp eq ptr null, %{{.*}} + // CHECK: icmp eq ptr %{{.*}}, null } void func(queue_t q); @@ -13,6 +13,6 @@ void func(queue_t q); void init(void) { queue_t q = 0; func(0); - // CHECK: store %opencl.queue_t* null, %opencl.queue_t** %q - // CHECK: call void @func(%opencl.queue_t* null) + // CHECK: store ptr null, ptr %q + // CHECK: call void @func(ptr null) } diff --git a/clang/test/CodeGenOpenCL/numbered-address-space.cl b/clang/test/CodeGenOpenCL/numbered-address-space.cl index 565c67b..13f8133 100644 --- a/clang/test/CodeGenOpenCL/numbered-address-space.cl +++ b/clang/test/CodeGenOpenCL/numbered-address-space.cl @@ -1,25 +1,24 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -no-opaque-pointers -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu tonga -S -emit-llvm -O0 -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu tonga -S -emit-llvm -O0 -o - %s | FileCheck %s // Make sure using numbered address spaces doesn't trigger crashes when a // builtin has an address space parameter. // CHECK-LABEL: @test_numbered_as_to_generic( -// CHECK: addrspacecast i32 addrspace(42)* %0 to i32* +// CHECK: addrspacecast ptr addrspace(42) %0 to ptr void test_numbered_as_to_generic(__attribute__((address_space(42))) int *arbitary_numbered_ptr) { generic int* generic_ptr = arbitary_numbered_ptr; *generic_ptr = 4; } // CHECK-LABEL: @test_generic_as_to_builtin_parameter_explicit_cast( -// CHECK: addrspacecast i32 addrspace(3)* %0 to i32* +// CHECK: addrspacecast ptr addrspace(3) %0 to ptr void test_generic_as_to_builtin_parameter_explicit_cast(__local int *local_ptr, float src) { generic int* generic_ptr = local_ptr; volatile float result = __builtin_amdgcn_ds_fmaxf((__local float*) generic_ptr, src, 0, 0, false); } // CHECK-LABEL: @test_generic_as_to_builtin_parameter_implicit_cast( -// CHECK: bitcast i32 addrspace(3)* %0 to float addrspace(3)* void test_generic_as_to_builtin_parameter_implicit_cast(__local int *local_ptr, float src) { volatile float result = __builtin_amdgcn_ds_fmaxf(local_ptr, src, 0, 0, false); } diff --git a/clang/test/CodeGenOpenCL/opencl_types.cl b/clang/test/CodeGenOpenCL/opencl_types.cl index 03f14b2..17e70a9 100644 --- a/clang/test/CodeGenOpenCL/opencl_types.cl +++ b/clang/test/CodeGenOpenCL/opencl_types.cl @@ -1,5 +1,5 @@ -// RUN: %clang_cc1 -no-opaque-pointers -cl-std=CL2.0 %s -triple "spir-unknown-unknown" -emit-llvm -o - -O0 | FileCheck %s --check-prefixes=CHECK-COM,CHECK-SPIR -// RUN: %clang_cc1 -no-opaque-pointers -cl-std=CL2.0 %s -triple "amdgcn--amdhsa" -emit-llvm -o - -O0 | FileCheck %s --check-prefixes=CHECK-COM,CHECK-AMDGCN +// RUN: %clang_cc1 -cl-std=CL2.0 %s -triple "spir-unknown-unknown" -emit-llvm -o - -O0 | FileCheck %s --check-prefixes=CHECK-COM,CHECK-SPIR +// RUN: %clang_cc1 -cl-std=CL2.0 %s -triple "amdgcn--amdhsa" -emit-llvm -o - -O0 | FileCheck %s --check-prefixes=CHECK-COM,CHECK-AMDGCN #define CLK_ADDRESS_CLAMP_TO_EDGE 2 #define CLK_NORMALIZED_COORDS_TRUE 1 @@ -10,67 +10,67 @@ constant sampler_t glb_smp = CLK_ADDRESS_CLAMP_TO_EDGE|CLK_NORMALIZED_COORDS_TRU // CHECK-COM-NOT: constant i32 void fnc1(image1d_t img) {} -// CHECK-SPIR: @fnc1(%opencl.image1d_ro_t addrspace(1)* -// CHECK-AMDGCN: @fnc1(%opencl.image1d_ro_t addrspace(4)* +// CHECK-SPIR: @fnc1(ptr addrspace(1) +// CHECK-AMDGCN: @fnc1(ptr addrspace(4) void fnc1arr(image1d_array_t img) {} -// CHECK-SPIR: @fnc1arr(%opencl.image1d_array_ro_t addrspace(1)* -// CHECK-AMDGCN: @fnc1arr(%opencl.image1d_array_ro_t addrspace(4)* +// CHECK-SPIR: @fnc1arr(ptr addrspace(1) +// CHECK-AMDGCN: @fnc1arr(ptr addrspace(4) void fnc1buff(image1d_buffer_t img) {} -// CHECK-SPIR: @fnc1buff(%opencl.image1d_buffer_ro_t addrspace(1)* -// CHECK-AMDGCN: @fnc1buff(%opencl.image1d_buffer_ro_t addrspace(4)* +// CHECK-SPIR: @fnc1buff(ptr addrspace(1) +// CHECK-AMDGCN: @fnc1buff(ptr addrspace(4) void fnc2(image2d_t img) {} -// CHECK-SPIR: @fnc2(%opencl.image2d_ro_t addrspace(1)* -// CHECK-AMDGCN: @fnc2(%opencl.image2d_ro_t addrspace(4)* +// CHECK-SPIR: @fnc2(ptr addrspace(1) +// CHECK-AMDGCN: @fnc2(ptr addrspace(4) void fnc2arr(image2d_array_t img) {} -// CHECK-SPIR: @fnc2arr(%opencl.image2d_array_ro_t addrspace(1)* -// CHECK-AMDGCN: @fnc2arr(%opencl.image2d_array_ro_t addrspace(4)* +// CHECK-SPIR: @fnc2arr(ptr addrspace(1) +// CHECK-AMDGCN: @fnc2arr(ptr addrspace(4) void fnc3(image3d_t img) {} -// CHECK-SPIR: @fnc3(%opencl.image3d_ro_t addrspace(1)* -// CHECK-AMDGCN: @fnc3(%opencl.image3d_ro_t addrspace(4)* +// CHECK-SPIR: @fnc3(ptr addrspace(1) +// CHECK-AMDGCN: @fnc3(ptr addrspace(4) void fnc4smp(sampler_t s) {} -// CHECK-SPIR-LABEL: define {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(2)* -// CHECK-AMDGCN-LABEL: define {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(4)* +// CHECK-SPIR-LABEL: define {{.*}}void @fnc4smp(ptr addrspace(2) +// CHECK-AMDGCN-LABEL: define {{.*}}void @fnc4smp(ptr addrspace(4) kernel void foo(image1d_t img) { sampler_t smp = CLK_ADDRESS_CLAMP_TO_EDGE|CLK_NORMALIZED_COORDS_TRUE|CLK_FILTER_LINEAR; - // CHECK-SPIR: alloca %opencl.sampler_t addrspace(2)* - // CHECK-AMDGCN: alloca %opencl.sampler_t addrspace(4)* + // CHECK-SPIR: alloca ptr addrspace(2) + // CHECK-AMDGCN: alloca ptr addrspace(4) event_t evt; - // CHECK-SPIR: alloca %opencl.event_t* - // CHECK-AMDGCN: alloca %opencl.event_t addrspace(5)* + // CHECK-SPIR: alloca ptr + // CHECK-AMDGCN: alloca ptr addrspace(5) clk_event_t clk_evt; - // CHECK-SPIR: alloca %opencl.clk_event_t* - // CHECK-AMDGCN: alloca %opencl.clk_event_t addrspace(1)* + // CHECK-SPIR: alloca ptr + // CHECK-AMDGCN: alloca ptr addrspace(1) queue_t queue; - // CHECK-SPIR: alloca %opencl.queue_t* - // CHECK-AMDGCN: alloca %opencl.queue_t addrspace(1)* + // CHECK-SPIR: alloca ptr + // CHECK-AMDGCN: alloca ptr addrspace(1) reserve_id_t rid; - // CHECK-SPIR: alloca %opencl.reserve_id_t* - // CHECK-AMDGCN: alloca %opencl.reserve_id_t addrspace(1)* - // CHECK-SPIR: store %opencl.sampler_t addrspace(2)* - // CHECK-AMDGCN: store %opencl.sampler_t addrspace(4)* + // CHECK-SPIR: alloca ptr + // CHECK-AMDGCN: alloca ptr addrspace(1) + // CHECK-SPIR: store ptr addrspace(2) + // CHECK-AMDGCN: store ptr addrspace(4) fnc4smp(smp); - // CHECK-SPIR: call {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(2)* - // CHECK-AMDGCN: call {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(4)* + // CHECK-SPIR: call {{.*}}void @fnc4smp(ptr addrspace(2) + // CHECK-AMDGCN: call {{.*}}void @fnc4smp(ptr addrspace(4) fnc4smp(glb_smp); - // CHECK-SPIR: call {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(2)* - // CHECK-AMDGCN: call {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(4)* + // CHECK-SPIR: call {{.*}}void @fnc4smp(ptr addrspace(2) + // CHECK-AMDGCN: call {{.*}}void @fnc4smp(ptr addrspace(4) } kernel void foo_ro_pipe(read_only pipe int p) {} -// CHECK-SPIR: @foo_ro_pipe(%opencl.pipe_ro_t addrspace(1)* %p) -// CHECK_AMDGCN: @foo_ro_pipe(%opencl.pipe_ro_t addrspace(1)* %p) +// CHECK-SPIR: @foo_ro_pipe(ptr addrspace(1) %p) +// CHECK_AMDGCN: @foo_ro_pipe(ptr addrspace(1) %p) kernel void foo_wo_pipe(write_only pipe int p) {} -// CHECK-SPIR: @foo_wo_pipe(%opencl.pipe_wo_t addrspace(1)* %p) -// CHECK_AMDGCN: @foo_wo_pipe(%opencl.pipe_wo_t addrspace(1)* %p) +// CHECK-SPIR: @foo_wo_pipe(ptr addrspace(1) %p) +// CHECK_AMDGCN: @foo_wo_pipe(ptr addrspace(1) %p) void __attribute__((overloadable)) bad1(image1d_t b, image2d_t c, image2d_t d) {} // CHECK-SPIR-LABEL: @{{_Z4bad114ocl_image1d_ro14ocl_image2d_roS0_|"\\01\?bad1@@\$\$J0YAXPAUocl_image1d_ro@@PAUocl_image2d_ro@@1@Z"}} -// CHECK-AMDGCN-LABEL: @{{_Z4bad114ocl_image1d_ro14ocl_image2d_roS0_|"\\01\?bad1@@\$\$J0YAXPAUocl_image1d_ro@@PAUocl_image2d_ro@@1@Z"}}(%opencl.image1d_ro_t addrspace(4)*{{.*}}%opencl.image2d_ro_t addrspace(4)*{{.*}}%opencl.image2d_ro_t addrspace(4)*{{.*}}) +// CHECK-AMDGCN-LABEL: @{{_Z4bad114ocl_image1d_ro14ocl_image2d_roS0_|"\\01\?bad1@@\$\$J0YAXPAUocl_image1d_ro@@PAUocl_image2d_ro@@1@Z"}}(ptr addrspace(4){{.*}}ptr addrspace(4){{.*}}ptr addrspace(4){{.*}}) diff --git a/clang/test/CodeGenOpenCL/overload.cl b/clang/test/CodeGenOpenCL/overload.cl index 62ed84c..589ab0a 100644 --- a/clang/test/CodeGenOpenCL/overload.cl +++ b/clang/test/CodeGenOpenCL/overload.cl @@ -1,5 +1,5 @@ -// RUN: %clang_cc1 -no-opaque-pointers -cl-std=CL2.0 -emit-llvm -o - -triple spir-unknown-unknown %s | FileCheck %s -// RUN: %clang_cc1 -no-opaque-pointers -cl-std=CL3.0 -cl-ext=+__opencl_c_generic_address_space -emit-llvm -o - -triple spir-unknown-unknown %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -emit-llvm -o - -triple spir-unknown-unknown %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL3.0 -cl-ext=+__opencl_c_generic_address_space -emit-llvm -o - -triple spir-unknown-unknown %s | FileCheck %s typedef short short4 __attribute__((ext_vector_type(4))); @@ -21,18 +21,18 @@ void kernel test1() { generic int *generic *gengen = 0; generic int *local *genloc = 0; generic int *global *genglob = 0; - // CHECK-DAG: call spir_func void @_Z3fooPU3AS1iS0_(i32 addrspace(1)* noundef {{.*}}, i32 addrspace(1)* noundef {{.*}}) + // CHECK-DAG: call spir_func void @_Z3fooPU3AS1iS0_(ptr addrspace(1) noundef {{.*}}, ptr addrspace(1) noundef {{.*}}) foo(a, b); - // CHECK-DAG: call spir_func void @_Z3fooPU3AS4iS0_(i32 addrspace(4)* noundef {{.*}}, i32 addrspace(4)* noundef {{.*}}) + // CHECK-DAG: call spir_func void @_Z3fooPU3AS4iS0_(ptr addrspace(4) noundef {{.*}}, ptr addrspace(4) noundef {{.*}}) foo(b, c); - // CHECK-DAG: call spir_func void @_Z3fooPU3AS4iS0_(i32 addrspace(4)* noundef {{.*}}, i32 addrspace(4)* noundef {{.*}}) + // CHECK-DAG: call spir_func void @_Z3fooPU3AS4iS0_(ptr addrspace(4) noundef {{.*}}, ptr addrspace(4) noundef {{.*}}) foo(a, d); - // CHECK-DAG: call spir_func void @_Z3barPU3AS4PU3AS4iS2_(i32 addrspace(4)* addrspace(4)* noundef {{.*}}, i32 addrspace(4)* addrspace(4)* noundef {{.*}}) + // CHECK-DAG: call spir_func void @_Z3barPU3AS4PU3AS4iS2_(ptr addrspace(4) noundef {{.*}}, ptr addrspace(4) noundef {{.*}}) bar(gengen, genloc); - // CHECK-DAG: call spir_func void @_Z3barPU3AS4PU3AS4iS2_(i32 addrspace(4)* addrspace(4)* noundef {{.*}}, i32 addrspace(4)* addrspace(4)* noundef {{.*}}) + // CHECK-DAG: call spir_func void @_Z3barPU3AS4PU3AS4iS2_(ptr addrspace(4) noundef {{.*}}, ptr addrspace(4) noundef {{.*}}) bar(gengen, genglob); - // CHECK-DAG: call spir_func void @_Z3barPU3AS1PU3AS4iS2_(i32 addrspace(4)* addrspace(1)* noundef {{.*}}, i32 addrspace(4)* addrspace(1)* noundef {{.*}}) + // CHECK-DAG: call spir_func void @_Z3barPU3AS1PU3AS4iS2_(ptr addrspace(1) noundef {{.*}}, ptr addrspace(1) noundef {{.*}}) bar(genglob, genglob); } diff --git a/clang/test/CodeGenOpenCL/partial_initializer.cl b/clang/test/CodeGenOpenCL/partial_initializer.cl index 483bdaf..5e246ca 100644 --- a/clang/test/CodeGenOpenCL/partial_initializer.cl +++ b/clang/test/CodeGenOpenCL/partial_initializer.cl @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -no-opaque-pointers -triple spir-unknown-unknown -cl-std=CL2.0 -emit-llvm %s -O0 -o - | FileCheck %s +// RUN: %clang_cc1 -triple spir-unknown-unknown -cl-std=CL2.0 -emit-llvm %s -O0 -o - | FileCheck %s typedef __attribute__(( ext_vector_type(2) )) int int2; typedef __attribute__(( ext_vector_type(4) )) int int4; @@ -35,32 +35,29 @@ void f(void) { // CHECK: %[[compoundliteral1:.*]] = alloca <2 x i32>, align 8 // CHECK: %[[V2:.*]] = alloca <4 x i32>, align 16 - // CHECK: %[[v0:.*]] = bitcast [6 x [6 x float]]* %A to i8* - // CHECK: call void @llvm.memset.p0i8.i32(i8* align 4 %[[v0]], i8 0, i32 144, i1 false) - // CHECK: %[[v1:.*]] = bitcast i8* %[[v0]] to [6 x [6 x float]]* - // CHECK: %[[v2:.*]] = getelementptr inbounds [6 x [6 x float]], [6 x [6 x float]]* %[[v1]], i32 0, i32 0 - // CHECK: %[[v3:.*]] = getelementptr inbounds [6 x float], [6 x float]* %[[v2]], i32 0, i32 0 - // CHECK: store float 1.000000e+00, float* %[[v3]], align 4 - // CHECK: %[[v4:.*]] = getelementptr inbounds [6 x float], [6 x float]* %[[v2]], i32 0, i32 1 - // CHECK: store float 2.000000e+00, float* %[[v4]], align 4 + // CHECK: call void @llvm.memset.p0.i32(ptr align 4 %A, i8 0, i32 144, i1 false) + // CHECK: %[[v2:.*]] = getelementptr inbounds [6 x [6 x float]], ptr %A, i32 0, i32 0 + // CHECK: %[[v3:.*]] = getelementptr inbounds [6 x float], ptr %[[v2]], i32 0, i32 0 + // CHECK: store float 1.000000e+00, ptr %[[v3]], align 4 + // CHECK: %[[v4:.*]] = getelementptr inbounds [6 x float], ptr %[[v2]], i32 0, i32 1 + // CHECK: store float 2.000000e+00, ptr %[[v4]], align 4 float A[6][6] = {1.0f, 2.0f}; - // CHECK: %[[v5:.*]] = bitcast %struct.StrucTy* %S to i8* - // CHECK: call void @llvm.memcpy.p0i8.p2i8.i32(i8* align 4 %[[v5]], i8 addrspace(2)* align 4 bitcast (%struct.StrucTy addrspace(2)* @__const.f.S to i8 addrspace(2)*), i32 12, i1 false) + // CHECK: call void @llvm.memcpy.p0.p2.i32(ptr align 4 %S, ptr addrspace(2) align 4 @__const.f.S, i32 12, i1 false) StrucTy S = {1, 2}; - // CHECK: store <2 x i32> <i32 1, i32 2>, <2 x i32>* %[[compoundliteral1]], align 8 - // CHECK: %[[v6:.*]] = load <2 x i32>, <2 x i32>* %[[compoundliteral1]], align 8 + // CHECK: store <2 x i32> <i32 1, i32 2>, ptr %[[compoundliteral1]], align 8 + // CHECK: %[[v6:.*]] = load <2 x i32>, ptr %[[compoundliteral1]], align 8 // CHECK: %[[vext:.*]] = shufflevector <2 x i32> %[[v6]], <2 x i32> poison, <4 x i32> <i32 0, i32 1, i32 undef, i32 undef> // CHECK: %[[vecinit:.*]] = shufflevector <4 x i32> %[[vext]], <4 x i32> undef, <4 x i32> <i32 0, i32 1, i32 undef, i32 undef> // CHECK: %[[vecinit2:.*]] = insertelement <4 x i32> %[[vecinit]], i32 3, i32 2 // CHECK: %[[vecinit3:.*]] = insertelement <4 x i32> %[[vecinit2]], i32 4, i32 3 - // CHECK: store <4 x i32> %[[vecinit3]], <4 x i32>* %[[compoundliteral]], align 16 - // CHECK: %[[v7:.*]] = load <4 x i32>, <4 x i32>* %[[compoundliteral]], align 16 - // CHECK: store <4 x i32> %[[v7]], <4 x i32>* %[[V1]], align 16 + // CHECK: store <4 x i32> %[[vecinit3]], ptr %[[compoundliteral]], align 16 + // CHECK: %[[v7:.*]] = load <4 x i32>, ptr %[[compoundliteral]], align 16 + // CHECK: store <4 x i32> %[[v7]], ptr %[[V1]], align 16 int4 V1 = (int4)((int2)(1,2),3,4); - // CHECK: store <4 x i32> <i32 1, i32 1, i32 1, i32 1>, <4 x i32>* %[[V2]], align 16 + // CHECK: store <4 x i32> <i32 1, i32 1, i32 1, i32 1>, ptr %[[V2]], align 16 int4 V2 = (int4)(1); } diff --git a/clang/test/CodeGenOpenCL/preserve_vec3.cl b/clang/test/CodeGenOpenCL/preserve_vec3.cl index 29fd25c..d91df75 100644 --- a/clang/test/CodeGenOpenCL/preserve_vec3.cl +++ b/clang/test/CodeGenOpenCL/preserve_vec3.cl @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -no-opaque-pointers %s -emit-llvm -o - -triple spir-unknown-unknown -fpreserve-vec3-type | FileCheck %s +// RUN: %clang_cc1 %s -emit-llvm -o - -triple spir-unknown-unknown -fpreserve-vec3-type | FileCheck %s typedef char char3 __attribute__((ext_vector_type(3))); typedef char char8 __attribute__((ext_vector_type(8))); @@ -9,58 +9,54 @@ typedef float float4 __attribute__((ext_vector_type(4))); void kernel foo(global float3 *a, global float3 *b) { // CHECK-LABEL: spir_kernel void @foo - // 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 + // CHECK: %[[LOAD_A:.*]] = load <3 x float>, ptr addrspace(1) %a + // CHECK: store <3 x float> %[[LOAD_A]], ptr addrspace(1) %b *b = *a; } void kernel float4_to_float3(global float3 *a, global float4 *b) { // CHECK-LABEL: spir_kernel void @float4_to_float3 - // CHECK: %[[LOAD_A:.*]] = load <4 x float>, <4 x float> addrspace(1)* %b, align 16 + // CHECK: %[[LOAD_A:.*]] = load <4 x float>, ptr addrspace(1) %b, align 16 // CHECK: %[[ASTYPE:.*]] = shufflevector <4 x float> %[[LOAD_A]], <4 x float> poison, <3 x i32> <i32 0, i32 1, i32 2> - // CHECK: store <3 x float> %[[ASTYPE]], <3 x float> addrspace(1)* %a, align 16 + // CHECK: store <3 x float> %[[ASTYPE]], ptr addrspace(1) %a, align 16 *a = __builtin_astype(*b, float3); } void kernel float3_to_float4(global float3 *a, global float4 *b) { // CHECK-LABEL: spir_kernel void @float3_to_float4 - // CHECK: %[[LOAD_A:.*]] = load <3 x float>, <3 x float> addrspace(1)* %a, align 16 + // CHECK: %[[LOAD_A:.*]] = load <3 x float>, ptr addrspace(1) %a, align 16 // CHECK: %[[ASTYPE:.*]] = shufflevector <3 x float> %[[LOAD_A]], <3 x float> poison, <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 + // CHECK: store <4 x float> %[[ASTYPE]], ptr addrspace(1) %b, align 16 *b = __builtin_astype(*a, float4); } void kernel float3_to_double2(global float3 *a, global double2 *b) { // CHECK-LABEL: spir_kernel void @float3_to_double2 - // CHECK: %[[LOAD_A:.*]] = load <3 x float>, <3 x float> addrspace(1)* %a, align 16 + // CHECK: %[[LOAD_A:.*]] = load <3 x float>, ptr addrspace(1) %a, align 16 // CHECK: %[[ASTYPE:.*]] = shufflevector <3 x float> %[[LOAD_A]], <3 x float> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 undef> - // CHECK: %[[OUT_BC:.*]] = bitcast <2 x double> addrspace(1)* %b to <4 x float> addrspace(1)* - // CHECK: store <4 x float> %[[ASTYPE]], <4 x float> addrspace(1)* %[[OUT_BC]], align 16 + // CHECK: store <4 x float> %[[ASTYPE]], ptr addrspace(1) %b, align 16 *b = __builtin_astype(*a, double2); } void kernel char8_to_short3(global short3 *a, global char8 *b) { // CHECK-LABEL: spir_kernel void @char8_to_short3 - // CHECK: %[[IN_BC:.*]] = bitcast <8 x i8> addrspace(1)* %b to <4 x i16> addrspace(1)* - // CHECK: %[[LOAD_B:.*]] = load <4 x i16>, <4 x i16> addrspace(1)* %[[IN_BC]] + // CHECK: %[[LOAD_B:.*]] = load <4 x i16>, ptr addrspace(1) %b // CHECK: %[[ASTYPE:.*]] = shufflevector <4 x i16> %[[LOAD_B]], <4 x i16> poison, <3 x i32> <i32 0, i32 1, i32 2> - // CHECK: store <3 x i16> %[[ASTYPE]], <3 x i16> addrspace(1)* %a, align 8 + // CHECK: store <3 x i16> %[[ASTYPE]], ptr addrspace(1) %a, align 8 *a = __builtin_astype(*b, short3); } void from_char3(char3 a, global int *out) { // CHECK-LABEL: void @from_char3 // CHECK: %[[ASTYPE:.*]] = shufflevector <3 x i8> %a, <3 x i8> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 undef> - // CHECK: %[[OUT_BC:.*]] = bitcast i32 addrspace(1)* %out to <4 x i8> addrspace(1)* - // CHECK: store <4 x i8> %[[ASTYPE]], <4 x i8> addrspace(1)* %[[OUT_BC]] + // CHECK: store <4 x i8> %[[ASTYPE]], ptr addrspace(1) %out *out = __builtin_astype(a, int); } void from_short3(short3 a, global long *out) { // CHECK-LABEL: void @from_short3 // CHECK: %[[ASTYPE:.*]] = shufflevector <3 x i16> %a, <3 x i16> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 undef> - // CHECK: %[[OUT_BC:.*]] = bitcast i64 addrspace(1)* %out to <4 x i16> addrspace(1)* - // CHECK: store <4 x i16> %[[ASTYPE]], <4 x i16> addrspace(1)* %[[OUT_BC]] + // CHECK: store <4 x i16> %[[ASTYPE]], ptr addrspace(1) %out *out = __builtin_astype(a, long); } @@ -68,7 +64,7 @@ void scalar_to_char3(int a, global char3 *out) { // CHECK-LABEL: void @scalar_to_char3 // CHECK: %[[IN_BC:.*]] = bitcast i32 %a to <4 x i8> // CHECK: %[[ASTYPE:.*]] = shufflevector <4 x i8> %[[IN_BC]], <4 x i8> poison, <3 x i32> <i32 0, i32 1, i32 2> - // CHECK: store <3 x i8> %[[ASTYPE]], <3 x i8> addrspace(1)* %out + // CHECK: store <3 x i8> %[[ASTYPE]], ptr addrspace(1) %out *out = __builtin_astype(a, char3); } @@ -76,6 +72,6 @@ void scalar_to_short3(long a, global short3 *out) { // CHECK-LABEL: void @scalar_to_short3 // CHECK: %[[IN_BC:.*]] = bitcast i64 %a to <4 x i16> // CHECK: %[[ASTYPE:.*]] = shufflevector <4 x i16> %[[IN_BC]], <4 x i16> poison, <3 x i32> <i32 0, i32 1, i32 2> - // CHECK: store <3 x i16> %[[ASTYPE]], <3 x i16> addrspace(1)* %out + // CHECK: store <3 x i16> %[[ASTYPE]], ptr addrspace(1) %out *out = __builtin_astype(a, short3); } diff --git a/clang/test/CodeGenOpenCL/printf.cl b/clang/test/CodeGenOpenCL/printf.cl index 3b6bffb..2e11b88 100644 --- a/clang/test/CodeGenOpenCL/printf.cl +++ b/clang/test/CodeGenOpenCL/printf.cl @@ -1,9 +1,9 @@ -// RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -cl-std=CL1.2 -cl-ext=-+cl_khr_fp64 -triple spir-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck -check-prefixes=FP64,ALL %s -// RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -cl-std=CL1.2 -cl-ext=-cl_khr_fp64 -triple spir-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck -check-prefixes=NOFP64,ALL %s -// RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -cl-std=CL3.0 -cl-ext=+__opencl_c_fp64,+cl_khr_fp64 -triple spir-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck -check-prefixes=FP64,ALL %s -// RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -cl-std=CL3.0 -cl-ext=-__opencl_c_fp64,-cl_khr_fp64 -triple spir-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck -check-prefixes=NOFP64,ALL %s -// RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -cl-std=clc++2021 -cl-ext=+__opencl_c_fp64,+cl_khr_fp64 -triple spir-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck -check-prefixes=FP64,ALL %s -// RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -cl-std=clc++2021 -cl-ext=-__opencl_c_fp64,-cl_khr_fp64 -triple spir-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck -check-prefixes=NOFP64,ALL %s +// RUN: %clang_cc1 -no-enable-noundef-analysis -cl-std=CL1.2 -cl-ext=-+cl_khr_fp64 -triple spir-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck -check-prefixes=FP64,ALL %s +// RUN: %clang_cc1 -no-enable-noundef-analysis -cl-std=CL1.2 -cl-ext=-cl_khr_fp64 -triple spir-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck -check-prefixes=NOFP64,ALL %s +// RUN: %clang_cc1 -no-enable-noundef-analysis -cl-std=CL3.0 -cl-ext=+__opencl_c_fp64,+cl_khr_fp64 -triple spir-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck -check-prefixes=FP64,ALL %s +// RUN: %clang_cc1 -no-enable-noundef-analysis -cl-std=CL3.0 -cl-ext=-__opencl_c_fp64,-cl_khr_fp64 -triple spir-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck -check-prefixes=NOFP64,ALL %s +// RUN: %clang_cc1 -no-enable-noundef-analysis -cl-std=clc++2021 -cl-ext=+__opencl_c_fp64,+cl_khr_fp64 -triple spir-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck -check-prefixes=FP64,ALL %s +// RUN: %clang_cc1 -no-enable-noundef-analysis -cl-std=clc++2021 -cl-ext=-__opencl_c_fp64,-cl_khr_fp64 -triple spir-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck -check-prefixes=NOFP64,ALL %s typedef __attribute__((ext_vector_type(2))) float float2; typedef __attribute__((ext_vector_type(2))) half half2; @@ -16,24 +16,24 @@ int printf(__constant const char* st, ...) __attribute__((format(printf, 1, 2))) // ALL-LABEL: @test_printf_float2( -// FP64: %call = call spir_func i32 (i8 addrspace(2)*, ...) @{{.*}}printf{{.*}}(i8 addrspace(2)* getelementptr inbounds ([7 x i8], [7 x i8] addrspace(2)* @.str, i32 0, i32 0), <2 x float> %0) +// FP64: %call = call spir_func i32 (ptr addrspace(2), ...) @{{.*}}printf{{.*}}(ptr addrspace(2) @.str, <2 x float> %0) -// NOFP64: call spir_func i32 (i8 addrspace(2)*, ...) @{{.*}}printf{{.*}}(i8 addrspace(2)* getelementptr inbounds ([7 x i8], [7 x i8] addrspace(2)* @.str, i32 0, i32 0), <2 x float> %0) +// NOFP64: call spir_func i32 (ptr addrspace(2), ...) @{{.*}}printf{{.*}}(ptr addrspace(2) @.str, <2 x float> %0) kernel void test_printf_float2(float2 arg) { printf("%v2hlf", arg); } // ALL-LABEL: @test_printf_half2( -// FP64: %call = call spir_func i32 (i8 addrspace(2)*, ...) @{{.*}}printf{{.*}}(i8 addrspace(2)* getelementptr inbounds ([6 x i8], [6 x i8] addrspace(2)* @.str.1, i32 0, i32 0), <2 x half> %0) +// FP64: %call = call spir_func i32 (ptr addrspace(2), ...) @{{.*}}printf{{.*}}(ptr addrspace(2) @.str.1, <2 x half> %0) -// NOFP64: %call = call spir_func i32 (i8 addrspace(2)*, ...) @{{.*}}printf{{.*}}(i8 addrspace(2)* getelementptr inbounds ([6 x i8], [6 x i8] addrspace(2)* @.str.1, i32 0, i32 0), <2 x half> %0) +// NOFP64: %call = call spir_func i32 (ptr addrspace(2), ...) @{{.*}}printf{{.*}}(ptr addrspace(2) @.str.1, <2 x half> %0) kernel void test_printf_half2(half2 arg) { printf("%v2hf", arg); } #if defined(cl_khr_fp64) || defined(__opencl_c_fp64) // FP64-LABEL: @test_printf_double2( -// FP64: call spir_func i32 (i8 addrspace(2)*, ...) @{{.*}}printf{{.*}}(i8 addrspace(2)* getelementptr inbounds ([6 x i8], [6 x i8] addrspace(2)* @.str.2, i32 0, i32 0), <2 x double> %0) +// FP64: call spir_func i32 (ptr addrspace(2), ...) @{{.*}}printf{{.*}}(ptr addrspace(2) @.str.2, <2 x double> %0) kernel void test_printf_double2(double2 arg) { printf("%v2lf", arg); } diff --git a/clang/test/CodeGenOpenCL/private-array-initialization.cl b/clang/test/CodeGenOpenCL/private-array-initialization.cl index 9e7ef2d..11f2bc36 100644 --- a/clang/test/CodeGenOpenCL/private-array-initialization.cl +++ b/clang/test/CodeGenOpenCL/private-array-initialization.cl @@ -1,29 +1,23 @@ -// RUN: %clang_cc1 -no-opaque-pointers %s -triple spir-unknown-unknown -O0 -emit-llvm -o - | FileCheck -check-prefix=PRIVATE0 %s -// RUN: %clang_cc1 -no-opaque-pointers %s -triple amdgcn-amd-amdhsa-unknown -O0 -emit-llvm -o - | FileCheck -check-prefix=PRIVATE5 %s +// RUN: %clang_cc1 %s -triple spir-unknown-unknown -O0 -emit-llvm -o - | FileCheck -check-prefix=PRIVATE0 %s +// RUN: %clang_cc1 %s -triple amdgcn-amd-amdhsa-unknown -O0 -emit-llvm -o - | FileCheck -check-prefix=PRIVATE5 %s // CHECK: @test.arr = private unnamed_addr addrspace(2) constant [3 x i32] [i32 1, i32 2, i32 3], align 4 void test() { __private int arr[] = {1, 2, 3}; -// PRIVATE0: %[[arr_i8_ptr:[0-9]+]] = bitcast [3 x i32]* %arr to i8* -// PRIVATE0: call void @llvm.memcpy.p0i8.p2i8.i32(i8* align 4 %[[arr_i8_ptr]], i8 addrspace(2)* align 4 bitcast ([3 x i32] addrspace(2)* @__const.test.arr to i8 addrspace(2)*), i32 12, i1 false) +// PRIVATE0: call void @llvm.memcpy.p0.p2.i32(ptr align 4 %arr, ptr addrspace(2) align 4 @__const.test.arr, i32 12, i1 false) // PRIVATE5: %arr = alloca [3 x i32], align 4, addrspace(5) -// PRIVATE5: %0 = bitcast [3 x i32] addrspace(5)* %arr to i8 addrspace(5)* -// PRIVATE5: call void @llvm.memcpy.p5i8.p4i8.i64(i8 addrspace(5)* align 4 %0, i8 addrspace(4)* align 4 bitcast ([3 x i32] addrspace(4)* @__const.test.arr to i8 addrspace(4)*), i64 12, i1 false) +// PRIVATE5: call void @llvm.memcpy.p5.p4.i64(ptr addrspace(5) align 4 %arr, ptr addrspace(4) align 4 @__const.test.arr, i64 12, i1 false) } __kernel void initializer_cast_is_valid_crash() { // PRIVATE0: %v512 = alloca [64 x i8], align 1 -// PRIVATE0: %0 = bitcast [64 x i8]* %v512 to i8* -// PRIVATE0: call void @llvm.memset.p0i8.i32(i8* align 1 %0, i8 0, i32 64, i1 false) -// PRIVATE0: %1 = bitcast i8* %0 to [64 x i8]* +// PRIVATE0: call void @llvm.memset.p0.i32(ptr align 1 %v512, i8 0, i32 64, i1 false) // PRIVATE5: %v512 = alloca [64 x i8], align 1, addrspace(5) -// PRIVATE5: %0 = bitcast [64 x i8] addrspace(5)* %v512 to i8 addrspace(5)* -// PRIVATE5: call void @llvm.memset.p5i8.i64(i8 addrspace(5)* align 1 %0, i8 0, i64 64, i1 false) -// PRIVATE5: %1 = bitcast i8 addrspace(5)* %0 to [64 x i8] addrspace(5)* +// PRIVATE5: call void @llvm.memset.p5.i64(ptr addrspace(5) align 1 %v512, i8 0, i64 64, i1 false) unsigned char v512[64] = { 0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00, 0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00, diff --git a/clang/test/CodeGenOpenCL/ptx-calls.cl b/clang/test/CodeGenOpenCL/ptx-calls.cl index f46d204..0081152 100644 --- a/clang/test/CodeGenOpenCL/ptx-calls.cl +++ b/clang/test/CodeGenOpenCL/ptx-calls.cl @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -no-opaque-pointers %s -triple nvptx-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s void device_function() { } @@ -9,5 +9,5 @@ __kernel void kernel_function() { } // CHECK-LABEL: define{{.*}} spir_kernel void @kernel_function() // CHECK: call void @device_function() -// CHECK: !{{[0-9]+}} = !{void ()* @kernel_function, !"kernel", i32 1} +// CHECK: !{{[0-9]+}} = !{ptr @kernel_function, !"kernel", i32 1} diff --git a/clang/test/CodeGenOpenCL/ptx-kernels.cl b/clang/test/CodeGenOpenCL/ptx-kernels.cl index 02a9f86..210e568 100644 --- a/clang/test/CodeGenOpenCL/ptx-kernels.cl +++ b/clang/test/CodeGenOpenCL/ptx-kernels.cl @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -no-opaque-pointers %s -triple nvptx-unknown-unknown -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -emit-llvm -o - | FileCheck %s void device_function() { } @@ -8,4 +8,4 @@ __kernel void kernel_function() { } // CHECK-LABEL: define{{.*}} spir_kernel void @kernel_function() -// CHECK: !{{[0-9]+}} = !{void ()* @kernel_function, !"kernel", i32 1} +// CHECK: !{{[0-9]+}} = !{ptr @kernel_function, !"kernel", i32 1} diff --git a/clang/test/CodeGenOpenCL/size_t.cl b/clang/test/CodeGenOpenCL/size_t.cl index 27db0f6..af09fb7 100644 --- a/clang/test/CodeGenOpenCL/size_t.cl +++ b/clang/test/CodeGenOpenCL/size_t.cl @@ -1,123 +1,123 @@ -// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL2.0 -finclude-default-header -fdeclare-opencl-builtins -emit-llvm -O0 -triple spir-unknown-unknown -o - | FileCheck --check-prefix=SZ32 %s -// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL2.0 -finclude-default-header -fdeclare-opencl-builtins -emit-llvm -O0 -triple spir64-unknown-unknown -o - | FileCheck --check-prefix=SZ64 --check-prefix=SZ64ONLY %s -// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL2.0 -finclude-default-header -fdeclare-opencl-builtins -emit-llvm -O0 -triple amdgcn -o - | FileCheck --check-prefix=SZ64 --check-prefix=AMDGCN %s -// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL2.0 -finclude-default-header -fdeclare-opencl-builtins -emit-llvm -O0 -triple amdgcn---opencl -o - | FileCheck --check-prefix=SZ64 --check-prefix=AMDGCN %s +// RUN: %clang_cc1 %s -cl-std=CL2.0 -finclude-default-header -fdeclare-opencl-builtins -emit-llvm -O0 -triple spir-unknown-unknown -o - | FileCheck --check-prefix=SZ32 %s +// RUN: %clang_cc1 %s -cl-std=CL2.0 -finclude-default-header -fdeclare-opencl-builtins -emit-llvm -O0 -triple spir64-unknown-unknown -o - | FileCheck --check-prefix=SZ64 --check-prefix=SZ64ONLY %s +// RUN: %clang_cc1 %s -cl-std=CL2.0 -finclude-default-header -fdeclare-opencl-builtins -emit-llvm -O0 -triple amdgcn -o - | FileCheck --check-prefix=SZ64 --check-prefix=AMDGCN %s +// RUN: %clang_cc1 %s -cl-std=CL2.0 -finclude-default-header -fdeclare-opencl-builtins -emit-llvm -O0 -triple amdgcn---opencl -o - | FileCheck --check-prefix=SZ64 --check-prefix=AMDGCN %s -//SZ32: define{{.*}} i32 @test_ptrtoint_private(i8* noundef %x) -//SZ32: ptrtoint i8* %{{.*}} to i32 -//SZ64ONLY: define{{.*}} i64 @test_ptrtoint_private(i8* noundef %x) -//SZ64ONLY: ptrtoint i8* %{{.*}} to i64 -//AMDGCN: define{{.*}} i64 @test_ptrtoint_private(i8 addrspace(5)* noundef %x) -//AMDGCN: ptrtoint i8 addrspace(5)* %{{.*}} to i64 +//SZ32: define{{.*}} i32 @test_ptrtoint_private(ptr noundef %x) +//SZ32: ptrtoint ptr %{{.*}} to i32 +//SZ64ONLY: define{{.*}} i64 @test_ptrtoint_private(ptr noundef %x) +//SZ64ONLY: ptrtoint ptr %{{.*}} to i64 +//AMDGCN: define{{.*}} i64 @test_ptrtoint_private(ptr addrspace(5) noundef %x) +//AMDGCN: ptrtoint ptr addrspace(5) %{{.*}} to i64 size_t test_ptrtoint_private(private char* x) { return (size_t)x; } -//SZ32: define{{.*}} i32 @test_ptrtoint_global(i8 addrspace(1)* noundef %x) -//SZ32: ptrtoint i8 addrspace(1)* %{{.*}} to i32 -//SZ64: define{{.*}} i64 @test_ptrtoint_global(i8 addrspace(1)* noundef %x) -//SZ64: ptrtoint i8 addrspace(1)* %{{.*}} to i64 +//SZ32: define{{.*}} i32 @test_ptrtoint_global(ptr addrspace(1) noundef %x) +//SZ32: ptrtoint ptr addrspace(1) %{{.*}} to i32 +//SZ64: define{{.*}} i64 @test_ptrtoint_global(ptr addrspace(1) noundef %x) +//SZ64: ptrtoint ptr addrspace(1) %{{.*}} to i64 intptr_t test_ptrtoint_global(global char* x) { return (intptr_t)x; } -//SZ32: define{{.*}} i32 @test_ptrtoint_constant(i8 addrspace(2)* noundef %x) -//SZ32: ptrtoint i8 addrspace(2)* %{{.*}} to i32 -//SZ64ONLY: define{{.*}} i64 @test_ptrtoint_constant(i8 addrspace(2)* noundef %x) -//SZ64ONLY: ptrtoint i8 addrspace(2)* %{{.*}} to i64 -//AMDGCN: define{{.*}} i64 @test_ptrtoint_constant(i8 addrspace(4)* noundef %x) -//AMDGCN: ptrtoint i8 addrspace(4)* %{{.*}} to i64 +//SZ32: define{{.*}} i32 @test_ptrtoint_constant(ptr addrspace(2) noundef %x) +//SZ32: ptrtoint ptr addrspace(2) %{{.*}} to i32 +//SZ64ONLY: define{{.*}} i64 @test_ptrtoint_constant(ptr addrspace(2) noundef %x) +//SZ64ONLY: ptrtoint ptr addrspace(2) %{{.*}} to i64 +//AMDGCN: define{{.*}} i64 @test_ptrtoint_constant(ptr addrspace(4) noundef %x) +//AMDGCN: ptrtoint ptr addrspace(4) %{{.*}} to i64 uintptr_t test_ptrtoint_constant(constant char* x) { return (uintptr_t)x; } -//SZ32: define{{.*}} i32 @test_ptrtoint_local(i8 addrspace(3)* noundef %x) -//SZ32: ptrtoint i8 addrspace(3)* %{{.*}} to i32 -//SZ64: define{{.*}} i64 @test_ptrtoint_local(i8 addrspace(3)* noundef %x) -//SZ64: ptrtoint i8 addrspace(3)* %{{.*}} to i64 +//SZ32: define{{.*}} i32 @test_ptrtoint_local(ptr addrspace(3) noundef %x) +//SZ32: ptrtoint ptr addrspace(3) %{{.*}} to i32 +//SZ64: define{{.*}} i64 @test_ptrtoint_local(ptr addrspace(3) noundef %x) +//SZ64: ptrtoint ptr addrspace(3) %{{.*}} to i64 size_t test_ptrtoint_local(local char* x) { return (size_t)x; } -//SZ32: define{{.*}} i32 @test_ptrtoint_generic(i8 addrspace(4)* noundef %x) -//SZ32: ptrtoint i8 addrspace(4)* %{{.*}} to i32 -//SZ64ONLY: define{{.*}} i64 @test_ptrtoint_generic(i8 addrspace(4)* noundef %x) -//SZ64ONLY: ptrtoint i8 addrspace(4)* %{{.*}} to i64 -//AMDGCN: define{{.*}} i64 @test_ptrtoint_generic(i8* noundef %x) -//AMDGCN: ptrtoint i8* %{{.*}} to i64 +//SZ32: define{{.*}} i32 @test_ptrtoint_generic(ptr addrspace(4) noundef %x) +//SZ32: ptrtoint ptr addrspace(4) %{{.*}} to i32 +//SZ64ONLY: define{{.*}} i64 @test_ptrtoint_generic(ptr addrspace(4) noundef %x) +//SZ64ONLY: ptrtoint ptr addrspace(4) %{{.*}} to i64 +//AMDGCN: define{{.*}} i64 @test_ptrtoint_generic(ptr noundef %x) +//AMDGCN: ptrtoint ptr %{{.*}} to i64 size_t test_ptrtoint_generic(generic char* x) { return (size_t)x; } -//SZ32: define{{.*}} i8* @test_inttoptr_private(i32 noundef %x) -//SZ32: inttoptr i32 %{{.*}} to i8* -//SZ64ONLY: define{{.*}} i8* @test_inttoptr_private(i64 noundef %x) -//SZ64ONLY: inttoptr i64 %{{.*}} to i8* -//AMDGCN: define{{.*}} i8 addrspace(5)* @test_inttoptr_private(i64 noundef %x) +//SZ32: define{{.*}} ptr @test_inttoptr_private(i32 noundef %x) +//SZ32: inttoptr i32 %{{.*}} to ptr +//SZ64ONLY: define{{.*}} ptr @test_inttoptr_private(i64 noundef %x) +//SZ64ONLY: inttoptr i64 %{{.*}} to ptr +//AMDGCN: define{{.*}} ptr addrspace(5) @test_inttoptr_private(i64 noundef %x) //AMDGCN: trunc i64 %{{.*}} to i32 -//AMDGCN: inttoptr i32 %{{.*}} to i8 addrspace(5)* +//AMDGCN: inttoptr i32 %{{.*}} to ptr addrspace(5) private char* test_inttoptr_private(size_t x) { return (private char*)x; } -//SZ32: define{{.*}} i8 addrspace(1)* @test_inttoptr_global(i32 noundef %x) -//SZ32: inttoptr i32 %{{.*}} to i8 addrspace(1)* -//SZ64: define{{.*}} i8 addrspace(1)* @test_inttoptr_global(i64 noundef %x) -//SZ64: inttoptr i64 %{{.*}} to i8 addrspace(1)* +//SZ32: define{{.*}} ptr addrspace(1) @test_inttoptr_global(i32 noundef %x) +//SZ32: inttoptr i32 %{{.*}} to ptr addrspace(1) +//SZ64: define{{.*}} ptr addrspace(1) @test_inttoptr_global(i64 noundef %x) +//SZ64: inttoptr i64 %{{.*}} to ptr addrspace(1) global char* test_inttoptr_global(size_t x) { return (global char*)x; } -//SZ32: define{{.*}} i8 addrspace(3)* @test_add_local(i8 addrspace(3)* noundef %x, i32 noundef %y) -//SZ32: getelementptr inbounds i8, i8 addrspace(3)* %{{.*}}, i32 -//SZ64: define{{.*}} i8 addrspace(3)* @test_add_local(i8 addrspace(3)* noundef %x, i64 noundef %y) +//SZ32: define{{.*}} ptr addrspace(3) @test_add_local(ptr addrspace(3) noundef %x, i32 noundef %y) +//SZ32: getelementptr inbounds i8, ptr addrspace(3) %{{.*}}, i32 +//SZ64: define{{.*}} ptr addrspace(3) @test_add_local(ptr addrspace(3) noundef %x, i64 noundef %y) //AMDGCN: trunc i64 %{{.*}} to i32 -//AMDGCN: getelementptr inbounds i8, i8 addrspace(3)* %{{.*}}, i32 -//SZ64ONLY: getelementptr inbounds i8, i8 addrspace(3)* %{{.*}}, i64 +//AMDGCN: getelementptr inbounds i8, ptr addrspace(3) %{{.*}}, i32 +//SZ64ONLY: getelementptr inbounds i8, ptr addrspace(3) %{{.*}}, i64 local char* test_add_local(local char* x, ptrdiff_t y) { return x + y; } -//SZ32: define{{.*}} i8 addrspace(1)* @test_add_global(i8 addrspace(1)* noundef %x, i32 noundef %y) -//SZ32: getelementptr inbounds i8, i8 addrspace(1)* %{{.*}}, i32 -//SZ64: define{{.*}} i8 addrspace(1)* @test_add_global(i8 addrspace(1)* noundef %x, i64 noundef %y) -//SZ64: getelementptr inbounds i8, i8 addrspace(1)* %{{.*}}, i64 +//SZ32: define{{.*}} ptr addrspace(1) @test_add_global(ptr addrspace(1) noundef %x, i32 noundef %y) +//SZ32: getelementptr inbounds i8, ptr addrspace(1) %{{.*}}, i32 +//SZ64: define{{.*}} ptr addrspace(1) @test_add_global(ptr addrspace(1) noundef %x, i64 noundef %y) +//SZ64: getelementptr inbounds i8, ptr addrspace(1) %{{.*}}, i64 global char* test_add_global(global char* x, ptrdiff_t y) { return x + y; } -//SZ32: define{{.*}} i32 @test_sub_local(i8 addrspace(3)* noundef %x, i8 addrspace(3)* noundef %y) -//SZ32: ptrtoint i8 addrspace(3)* %{{.*}} to i32 -//SZ32: ptrtoint i8 addrspace(3)* %{{.*}} to i32 -//SZ64: define{{.*}} i64 @test_sub_local(i8 addrspace(3)* noundef %x, i8 addrspace(3)* noundef %y) -//SZ64: ptrtoint i8 addrspace(3)* %{{.*}} to i64 -//SZ64: ptrtoint i8 addrspace(3)* %{{.*}} to i64 +//SZ32: define{{.*}} i32 @test_sub_local(ptr addrspace(3) noundef %x, ptr addrspace(3) noundef %y) +//SZ32: ptrtoint ptr addrspace(3) %{{.*}} to i32 +//SZ32: ptrtoint ptr addrspace(3) %{{.*}} to i32 +//SZ64: define{{.*}} i64 @test_sub_local(ptr addrspace(3) noundef %x, ptr addrspace(3) noundef %y) +//SZ64: ptrtoint ptr addrspace(3) %{{.*}} to i64 +//SZ64: ptrtoint ptr addrspace(3) %{{.*}} to i64 ptrdiff_t test_sub_local(local char* x, local char *y) { return x - y; } -//SZ32: define{{.*}} i32 @test_sub_private(i8* noundef %x, i8* noundef %y) -//SZ32: ptrtoint i8* %{{.*}} to i32 -//SZ32: ptrtoint i8* %{{.*}} to i32 -//SZ64ONLY: define{{.*}} i64 @test_sub_private(i8* noundef %x, i8* noundef %y) -//SZ64ONLY: ptrtoint i8* %{{.*}} to i64 -//SZ64ONLY: ptrtoint i8* %{{.*}} to i64 -//AMDGCN: define{{.*}} i64 @test_sub_private(i8 addrspace(5)* noundef %x, i8 addrspace(5)* noundef %y) -//AMDGCN: ptrtoint i8 addrspace(5)* %{{.*}} to i64 -//AMDGCN: ptrtoint i8 addrspace(5)* %{{.*}} to i64 +//SZ32: define{{.*}} i32 @test_sub_private(ptr noundef %x, ptr noundef %y) +//SZ32: ptrtoint ptr %{{.*}} to i32 +//SZ32: ptrtoint ptr %{{.*}} to i32 +//SZ64ONLY: define{{.*}} i64 @test_sub_private(ptr noundef %x, ptr noundef %y) +//SZ64ONLY: ptrtoint ptr %{{.*}} to i64 +//SZ64ONLY: ptrtoint ptr %{{.*}} to i64 +//AMDGCN: define{{.*}} i64 @test_sub_private(ptr addrspace(5) noundef %x, ptr addrspace(5) noundef %y) +//AMDGCN: ptrtoint ptr addrspace(5) %{{.*}} to i64 +//AMDGCN: ptrtoint ptr addrspace(5) %{{.*}} to i64 ptrdiff_t test_sub_private(private char* x, private char *y) { return x - y; } -//SZ32: define{{.*}} i32 @test_sub_mix(i8* noundef %x, i8 addrspace(4)* noundef %y) -//SZ32: ptrtoint i8* %{{.*}} to i32 -//SZ32: ptrtoint i8 addrspace(4)* %{{.*}} to i32 -//SZ64ONLY: define{{.*}} i64 @test_sub_mix(i8* noundef %x, i8 addrspace(4)* noundef %y) -//SZ64ONLY: ptrtoint i8* %{{.*}} to i64 -//SZ64ONLY: ptrtoint i8 addrspace(4)* %{{.*}} to i64 -//AMDGCN: define{{.*}} i64 @test_sub_mix(i8 addrspace(5)* noundef %x, i8* noundef %y) -//AMDGCN: ptrtoint i8 addrspace(5)* %{{.*}} to i64 -//AMDGCN: ptrtoint i8* %{{.*}} to i64 +//SZ32: define{{.*}} i32 @test_sub_mix(ptr noundef %x, ptr addrspace(4) noundef %y) +//SZ32: ptrtoint ptr %{{.*}} to i32 +//SZ32: ptrtoint ptr addrspace(4) %{{.*}} to i32 +//SZ64ONLY: define{{.*}} i64 @test_sub_mix(ptr noundef %x, ptr addrspace(4) noundef %y) +//SZ64ONLY: ptrtoint ptr %{{.*}} to i64 +//SZ64ONLY: ptrtoint ptr addrspace(4) %{{.*}} to i64 +//AMDGCN: define{{.*}} i64 @test_sub_mix(ptr addrspace(5) noundef %x, ptr noundef %y) +//AMDGCN: ptrtoint ptr addrspace(5) %{{.*}} to i64 +//AMDGCN: ptrtoint ptr %{{.*}} to i64 ptrdiff_t test_sub_mix(private char* x, generic char *y) { return x - y; } diff --git a/clang/test/CodeGenOpenCL/spir-calling-conv.cl b/clang/test/CodeGenOpenCL/spir-calling-conv.cl index 8004a5c..569ea0c 100644 --- a/clang/test/CodeGenOpenCL/spir-calling-conv.cl +++ b/clang/test/CodeGenOpenCL/spir-calling-conv.cl @@ -1,18 +1,18 @@ -// RUN: %clang_cc1 -no-opaque-pointers %s -triple "spir-unknown-unknown" -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple "spir-unknown-unknown" -emit-llvm -o - | FileCheck %s int get_dummy_id(int D); kernel void bar(global int *A); kernel void foo(global int *A) -// CHECK: define{{.*}} spir_kernel void @foo(i32 addrspace(1)* noundef align 4 %A) +// CHECK: define{{.*}} spir_kernel void @foo(ptr addrspace(1) noundef align 4 %A) { int id = get_dummy_id(0); // CHECK: %{{[a-z0-9_]+}} = tail call spir_func i32 @get_dummy_id(i32 noundef 0) A[id] = id; bar(A); - // CHECK: tail call spir_kernel void @bar(i32 addrspace(1)* noundef align 4 %A) + // CHECK: tail call spir_kernel void @bar(ptr addrspace(1) noundef align 4 %A) } // CHECK: declare spir_func i32 @get_dummy_id(i32 noundef) -// CHECK: declare spir_kernel void @bar(i32 addrspace(1)* noundef align 4) +// CHECK: declare spir_kernel void @bar(ptr addrspace(1) noundef align 4) diff --git a/clang/test/CodeGenOpenCL/spir32_target.cl b/clang/test/CodeGenOpenCL/spir32_target.cl index f27e84c..924b2c1 100644 --- a/clang/test/CodeGenOpenCL/spir32_target.cl +++ b/clang/test/CodeGenOpenCL/spir32_target.cl @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -no-opaque-pointers %s -triple "spir-unknown-unknown" -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple "spir-unknown-unknown" -emit-llvm -o - | FileCheck %s // CHECK: target triple = "spir-unknown-unknown" @@ -16,7 +16,7 @@ kernel void foo(global long *arg) { my_st *tmp = 0; arg[0] = (long)(&tmp->v); -//CHECK: store i64 4, i64 addrspace(1)* +//CHECK: store i64 4, ptr addrspace(1) arg[1] = (long)(&tmp->v2); -//CHECK: store i64 8, i64 addrspace(1)* +//CHECK: store i64 8, ptr addrspace(1) } diff --git a/clang/test/CodeGenOpenCL/spir64_target.cl b/clang/test/CodeGenOpenCL/spir64_target.cl index c62dbb6..ba4a66d 100644 --- a/clang/test/CodeGenOpenCL/spir64_target.cl +++ b/clang/test/CodeGenOpenCL/spir64_target.cl @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -no-opaque-pointers %s -triple "spir64-unknown-unknown" -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple "spir64-unknown-unknown" -emit-llvm -o - | FileCheck %s // CHECK: target triple = "spir64-unknown-unknown" @@ -15,7 +15,7 @@ kernel void foo(global long *arg) { my_st *tmp = 0; arg[3] = (long)(&tmp->v); -//CHECK: store i64 8, i64 addrspace(1)* +//CHECK: store i64 8, ptr addrspace(1) arg[4] = (long)(&tmp->v2); -//CHECK: store i64 16, i64 addrspace(1)* +//CHECK: store i64 16, ptr addrspace(1) } diff --git a/clang/test/CodeGenOpenCL/spirv_target.cl b/clang/test/CodeGenOpenCL/spirv_target.cl index 3de62d0..2aeed19 100644 --- a/clang/test/CodeGenOpenCL/spirv_target.cl +++ b/clang/test/CodeGenOpenCL/spirv_target.cl @@ -1,5 +1,5 @@ -// RUN: %clang_cc1 -no-opaque-pointers %s -triple "spirv32-unknown-unknown" -verify -emit-llvm -o - | FileCheck %s -check-prefix=SPIRV32 -// RUN: %clang_cc1 -no-opaque-pointers %s -triple "spirv64-unknown-unknown" -verify -emit-llvm -o - | FileCheck %s -check-prefix=SPIRV64 +// RUN: %clang_cc1 %s -triple "spirv32-unknown-unknown" -verify -emit-llvm -o - | FileCheck %s -check-prefix=SPIRV32 +// RUN: %clang_cc1 %s -triple "spirv64-unknown-unknown" -verify -emit-llvm -o - | FileCheck %s -check-prefix=SPIRV64 // SPIRV32: target triple = "spirv32-unknown-unknown" // SPIRV64: target triple = "spirv64-unknown-unknown" @@ -22,10 +22,10 @@ kernel void foo(global long *arg) { #endif my_st *tmp = 0; - // SPIRV32: store i64 4, i64 addrspace(1)* - // SPIRV64: store i64 8, i64 addrspace(1)* + // SPIRV32: store i64 4, ptr addrspace(1) + // SPIRV64: store i64 8, ptr addrspace(1) arg[0] = (long)(&tmp->v); - // SPIRV32: store i64 8, i64 addrspace(1)* - // SPIRV64: store i64 16, i64 addrspace(1)* + // SPIRV32: store i64 8, ptr addrspace(1) + // SPIRV64: store i64 16, ptr addrspace(1) arg[1] = (long)(&tmp->v2); } diff --git a/clang/test/CodeGenOpenCL/str_literals.cl b/clang/test/CodeGenOpenCL/str_literals.cl index 0c25dbb..b1ebfc6 100644 --- a/clang/test/CodeGenOpenCL/str_literals.cl +++ b/clang/test/CodeGenOpenCL/str_literals.cl @@ -1,15 +1,15 @@ -// RUN: %clang_cc1 -no-opaque-pointers %s -cl-opt-disable -emit-llvm -o - -ffake-address-space-map | FileCheck %s +// RUN: %clang_cc1 %s -cl-opt-disable -emit-llvm -o - -ffake-address-space-map | FileCheck %s __constant char *__constant x = "hello world"; __constant char *__constant y = "hello world"; // CHECK: unnamed_addr addrspace(2) constant{{.*}}"hello world\00" // CHECK-NOT: addrspace(2) unnamed_addr constant -// CHECK: @x = {{(dso_local )?}}addrspace(2) constant i8 addrspace(2)* -// CHECK: @y = {{(dso_local )?}}addrspace(2) constant i8 addrspace(2)* +// CHECK: @x = {{(dso_local )?}}addrspace(2) constant ptr addrspace(2) +// CHECK: @y = {{(dso_local )?}}addrspace(2) constant ptr addrspace(2) // CHECK: unnamed_addr addrspace(2) constant{{.*}}"f\00" void f() { - //CHECK: store i8 addrspace(2)* {{.*}}, i8 addrspace(2)** + //CHECK: store ptr addrspace(2) {{.*}}, ptr constant const char *f3 = __func__; } diff --git a/clang/test/CodeGenOpenCL/vectorLoadStore.cl b/clang/test/CodeGenOpenCL/vectorLoadStore.cl index 15338b1..7f89929 100644 --- a/clang/test/CodeGenOpenCL/vectorLoadStore.cl +++ b/clang/test/CodeGenOpenCL/vectorLoadStore.cl @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -no-opaque-pointers -cl-std=CL2.0 -triple "spir-unknown-unknown" %s -emit-llvm -O0 -o - | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -triple "spir-unknown-unknown" %s -emit-llvm -O0 -o - | FileCheck %s typedef char char2 __attribute((ext_vector_type(2))); typedef char char3 __attribute((ext_vector_type(3))); @@ -16,7 +16,7 @@ void alignment() { __private char2 data_generic[100]; __private char8 data_private[100]; - // CHECK: %{{.*}} = load <4 x float>, <4 x float> addrspace(4)* %{{.*}}, align 2 - // CHECK: store <4 x float> %{{.*}}, <4 x float>* %{{.*}}, align 8 + // CHECK: %{{.*}} = load <4 x float>, ptr addrspace(4) %{{.*}}, align 2 + // CHECK: store <4 x float> %{{.*}}, ptr %{{.*}}, align 8 ((private float4 *)data_private)[1] = ((float4 *)data_generic)[2]; } diff --git a/clang/test/CodeGenOpenCL/vector_literals.cl b/clang/test/CodeGenOpenCL/vector_literals.cl index 587c569..d8d36c7 100644 --- a/clang/test/CodeGenOpenCL/vector_literals.cl +++ b/clang/test/CodeGenOpenCL/vector_literals.cl @@ -1,5 +1,5 @@ -// RUN: %clang_cc1 -no-opaque-pointers -emit-llvm %s -o - -O0 | FileCheck %s -// RUN: %clang_cc1 -no-opaque-pointers -emit-llvm %s -o - -cl-std=clc++ -O0 | FileCheck %s +// RUN: %clang_cc1 -emit-llvm %s -o - -O0 | FileCheck %s +// RUN: %clang_cc1 -emit-llvm %s -o - -cl-std=clc++ -O0 | FileCheck %s typedef __attribute__((ext_vector_type(2))) int int2; typedef __attribute__((ext_vector_type(3))) int int3; @@ -18,55 +18,55 @@ void vector_literals_valid() { //CHECK: insertelement <4 x i32> %{{.+}}, i32 %{{.+}}, i32 3 int4 a_1_1_1_1 = (int4)(1, 2, c1.s2, c2.s3); - //CHECK: store <2 x i32> <i32 1, i32 2>, <2 x i32>* + //CHECK: store <2 x i32> <i32 1, i32 2>, ptr //CHECK: shufflevector <2 x i32> %{{[0-9]+}}, <2 x i32> poison, <4 x i32> <i32 0, i32 1, i32 undef, i32 undef> //CHECK: shufflevector <4 x i32> %{{.+}}, <4 x i32> undef, <4 x i32> <i32 0, i32 1, i32 undef, i32 undef> //CHECK: insertelement <4 x i32> %{{.+}}, i32 3, i32 2 //CHECK: insertelement <4 x i32> %{{.+}}, i32 4, i32 3 int4 a_2_1_1 = (int4)((int2)(1, 2), 3, 4); - //CHECK: store <2 x i32> <i32 2, i32 3>, <2 x i32>* + //CHECK: store <2 x i32> <i32 2, i32 3>, ptr //CHECK: shufflevector <2 x i32> %{{[0-9]+}}, <2 x i32> poison, <4 x i32> <i32 0, i32 1, i32 undef, i32 undef> //CHECK: shufflevector <4 x i32> <i32 1, i32 undef, i32 undef, i32 undef>, <4 x i32> %{{.+}}, <4 x i32> <i32 0, i32 4, i32 5, i32 undef> //CHECK: insertelement <4 x i32> %{{.+}}, i32 4, i32 3 int4 a_1_2_1 = (int4)(1, (int2)(2, 3), 4); - //CHECK: store <2 x i32> <i32 3, i32 4>, <2 x i32>* + //CHECK: store <2 x i32> <i32 3, i32 4>, ptr //CHECK: shufflevector <2 x i32> %{{[0-9]+}}, <2 x i32> poison, <4 x i32> <i32 0, i32 1, i32 undef, i32 undef> //CHECK: shufflevector <4 x i32> <i32 1, i32 2, i32 undef, i32 undef>, <4 x i32> %{{.+}}, <4 x i32> <i32 0, i32 1, i32 4, i32 5> int4 a_1_1_2 = (int4)(1, 2, (int2)(3, 4)); - //CHECK: store <2 x i32> <i32 1, i32 2>, <2 x i32>* + //CHECK: store <2 x i32> <i32 1, i32 2>, ptr //CHECK: shufflevector <2 x i32> %{{[0-9]+}}, <2 x i32> poison, <4 x i32> <i32 0, i32 1, i32 undef, i32 undef> //CHECK: shufflevector <4 x i32> %{{.+}}, <4 x i32> undef, <4 x i32> <i32 0, i32 1, i32 undef, i32 undef> //CHECK: shufflevector <4 x i32> %{{.+}}, <4 x i32> <i32 3, i32 3, i32 undef, i32 undef>, <4 x i32> <i32 0, i32 1, i32 4, i32 5> int4 a_2_2 = (int4)((int2)(1, 2), (int2)(3)); - //CHECK: store <4 x i32> <i32 2, i32 3, i32 4, i32 undef>, <4 x i32>* + //CHECK: store <4 x i32> <i32 2, i32 3, i32 4, i32 undef>, ptr //CHECK: shufflevector <4 x i32> %{{.+}}, <4 x i32> poison, <3 x i32> <i32 0, i32 1, i32 2> //CHECK: shufflevector <3 x i32> %{{.+}}, <3 x i32> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 undef> //CHECK: shufflevector <4 x i32> <i32 1, i32 undef, i32 undef, i32 undef>, <4 x i32> %{{.+}}, <4 x i32> <i32 0, i32 4, i32 5, i32 6> int4 a_1_3 = (int4)(1, (int3)(2, 3, 4)); - //CHECK: store <4 x i32> <i32 1, i32 1, i32 1, i32 1>, <4 x i32>* %a + //CHECK: store <4 x i32> <i32 1, i32 1, i32 1, i32 1>, ptr %a int4 a = (int4)(1); - //CHECK: load <4 x i32>, <4 x i32>* %a + //CHECK: load <4 x i32>, ptr %a //CHECK: shufflevector <4 x i32> %{{[0-9]+}}, <4 x i32> poison, <2 x i32> <i32 0, i32 1> //CHECK: shufflevector <2 x i32> %{{[0-9]+}}, <2 x i32> poison, <8 x i32> <i32 0, i32 1, i32 undef, i32 undef, i32 undef, i32 undef, i32 undef, i32 undef> //CHECK: shufflevector <8 x i32> <i32 1, i32 2, i32 undef, i32 undef, i32 undef, i32 undef, i32 undef, i32 undef>, <8 x i32> %{{.+}}, <8 x i32> <i32 0, i32 1, i32 8, i32 9, i32 undef, i32 undef, i32 undef, i32 undef> - //CHECK: load <4 x i32>, <4 x i32>* %a + //CHECK: load <4 x i32>, ptr %a //CHECK: shufflevector <4 x i32> %{{[0-9]+}}, <4 x i32> poison, <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 undef, i32 undef, i32 undef, i32 undef> //CHECK: shufflevector <8 x i32> %{{.+}}, <8 x i32> %{{.+}}, <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 8, i32 9, i32 10, i32 11> int8 b = (int8)(1, 2, a.xy, a); - //CHECK: store <4 x float> <float 1.000000e+00, float 1.000000e+00, float 1.000000e+00, float 1.000000e+00>, <4 x float>* %V2 + //CHECK: store <4 x float> <float 1.000000e+00, float 1.000000e+00, float 1.000000e+00, float 1.000000e+00>, ptr %V2 float4 V2 = (float4)(1); } void vector_literals_with_cast() { // CHECK-LABEL: vector_literals_with_cast - // CHECK: store <2 x i32> <i32 12, i32 34>, <2 x i32>* + // CHECK: store <2 x i32> <i32 12, i32 34>, ptr // CHECK: extractelement <2 x i32> %{{[0-9]+}}, i{{[0-9]+}} 0 unsigned int withCast = ((int2)((int2)(12, 34))).s0; } |