[clang] aae20a7 - [CodeGenOpenCL] Convert some tests to opaque pointers (NFC)
Nikita Popov via cfe-commits
cfe-commits at lists.llvm.org
Thu Jan 5 01:58:47 PST 2023
Author: Nikita Popov
Date: 2023-01-05T10:57:30+01:00
New Revision: aae20a7421c5393316c25a8b614b370859c1a18f
URL: https://github.com/llvm/llvm-project/commit/aae20a7421c5393316c25a8b614b370859c1a18f
DIFF: https://github.com/llvm/llvm-project/commit/aae20a7421c5393316c25a8b614b370859c1a18f.diff
LOG: [CodeGenOpenCL] Convert some tests to opaque pointers (NFC)
Added:
Modified:
clang/test/CodeGenOpenCL/addr-space-struct-arg.cl
clang/test/CodeGenOpenCL/address-space-constant-initializers.cl
clang/test/CodeGenOpenCL/address-spaces.cl
clang/test/CodeGenOpenCL/amdgcn-automatic-variable.cl
clang/test/CodeGenOpenCL/amdgcn-non-temporal-store.cl
clang/test/CodeGenOpenCL/amdgpu-call-kernel.cl
clang/test/CodeGenOpenCL/amdgpu-debug-info-variable-expression.cl
clang/test/CodeGenOpenCL/amdgpu-printf.cl
clang/test/CodeGenOpenCL/as_type.cl
clang/test/CodeGenOpenCL/atomics-cas-remarks-gfx90a.cl
clang/test/CodeGenOpenCL/atomics-unsafe-hw-remarks-gfx90a.cl
clang/test/CodeGenOpenCL/bool_cast.cl
clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl
clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
clang/test/CodeGenOpenCL/builtins-amdgcn.cl
clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx8.cl
clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx90a.cl
clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx940.cl
clang/test/CodeGenOpenCL/builtins-generic-amdgcn.cl
clang/test/CodeGenOpenCL/builtins-r600.cl
clang/test/CodeGenOpenCL/byval.cl
clang/test/CodeGenOpenCL/constant-addr-space-globals.cl
clang/test/CodeGenOpenCL/event_t.cl
clang/test/CodeGenOpenCL/images.cl
clang/test/CodeGenOpenCL/kernel-param-alignment.cl
clang/test/CodeGenOpenCL/lifetime.cl
clang/test/CodeGenOpenCL/memcpy.cl
clang/test/CodeGenOpenCL/null_queue.cl
clang/test/CodeGenOpenCL/numbered-address-space.cl
clang/test/CodeGenOpenCL/opencl_types.cl
clang/test/CodeGenOpenCL/overload.cl
clang/test/CodeGenOpenCL/partial_initializer.cl
clang/test/CodeGenOpenCL/preserve_vec3.cl
clang/test/CodeGenOpenCL/printf.cl
clang/test/CodeGenOpenCL/private-array-initialization.cl
clang/test/CodeGenOpenCL/ptx-calls.cl
clang/test/CodeGenOpenCL/ptx-kernels.cl
clang/test/CodeGenOpenCL/size_t.cl
clang/test/CodeGenOpenCL/spir-calling-conv.cl
clang/test/CodeGenOpenCL/spir32_target.cl
clang/test/CodeGenOpenCL/spir64_target.cl
clang/test/CodeGenOpenCL/spirv_target.cl
clang/test/CodeGenOpenCL/str_literals.cl
clang/test/CodeGenOpenCL/vectorLoadStore.cl
clang/test/CodeGenOpenCL/vector_literals.cl
Removed:
################################################################################
diff --git a/clang/test/CodeGenOpenCL/addr-space-struct-arg.cl b/clang/test/CodeGenOpenCL/addr-space-struct-arg.cl
index 2f45dc2ede0f2..a806c5f1b6eb1 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 8b488d2391bae..14fbcba67ad7c 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 7ad8220c676e1..5b2a95c6ac16a 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 88754c7a35d7f..f26495bc44aa3 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 72f1a001ed65d..34938607a5e9c 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 22e0fd1c7f1aa..8ad0beac3dde5 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 425f764f2664a..ba645442f6a6f 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 1bc0e723e17a6..edf6dbf8657cb 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 355f02617d609..afc76b7b57a9d 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 c62d58f510bee..a5321ea7c158d 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 61a24acf31ac2..1243745c17b5b 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 9db3d4a270220..6e6538f00b4f2 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 0a752a44d4d69..da989ecba941f 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 c7437d7c1e881..590670b9786a5 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 9a20dba7490b4..82cd3177d6c6d 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 e4d85f71c63c3..f9782c16ab343 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 f078f4e93bb3d..c582c7fd91c27 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 8f22f829b6956..12b85937b9063 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 7ead1f0e80733..46c044a0bef69 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 274de1501ef45..7ceb8d68cae67 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 bd994ddca4e3c..6e734d7c5d83e 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 adb9252bdb45d..c38d3ed7ed7b7 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 b94332b286197..bb1041402b203 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 91867baad28f0..f15397a95bf18 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 d4e662825bfbc..c1a17407a2548 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 c016835d48af3..45df0c4038b05 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 5255a2ffc6a2a..06362fffadcd0 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 aaa94574bd08f..d6b3123a5a0b3 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 565c67bc6ff70..13f81330d4e7e 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 03f14b2b4475c..17e70a9f774fb 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 at Z"}}
-// CHECK-AMDGCN-LABEL: @{{_Z4bad114ocl_image1d_ro14ocl_image2d_roS0_|"\\01\?bad1@@\$\$J0YAXPAUocl_image1d_ro@@PAUocl_image2d_ro@@1 at 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 at 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 62ed84c7ec972..589ab0a6fd3c2 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 483bdaffa66a2..5e246ca71655a 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 29fd25ca288c8..d91df758546ed 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 3b6bffbf265e7..2e11b8889d23f 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 9e7ef2d6a5fbc..11f2bc362b849 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 f46d20461a805..0081152ae40e0 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 02a9f8647610e..210e5682ac721 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 27db0f64dd917..af09fb7d925f7 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, ptr
diff _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, ptr
diff _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
ptr
diff _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
ptr
diff _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
ptr
diff _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 8004a5cafc011..569ea0cbe1af6 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 f27e84c78cf23..924b2c12f5537 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 c62dbb6c1db22..ba4a66d7158fd 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 3de62d066e296..2aeed195449a8 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 0c25dbb907b24..b1ebfc65fdea7 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 15338b1cc55c1..7f8992929b489 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 587c56939a322..d8d36c7ea493f 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;
}
More information about the cfe-commits
mailing list