[clang] Convert a few tests to opaque pointers

[barannikov88]

convert a few more tests

[nikic]

LGTM

[barannikov88]

convert a few more tests, last iteration

[barannikov88]

@nikic I converted a few more tests after you approved the changes. Please take a look once more, or let me know if I should revert to the previous revision.

[nikic]

Still LGTM

[barannikov88]

Thanks. These were the lasts tests in these directories. The only one remaining is clang/test/CodeGenOpenCL/opaque-ptr-spirv.cl, I don't know what to do with it.

[amyk]

Hi, I just wanted to give a heads up that it appears that this patch is causing failures on some of the PPC buildbots.
In particular, the const-str-array-decay.cl appears to be failing. The failing build can be seen in: https://lab.llvm.org/buildbot/#/builders/36/builds/33110

[barannikov88]

In D150520#4342429, @amyk wrote:

Hi, I just wanted to give a heads up that it appears that this patch is causing failures on some of the PPC buildbots.
In particular, the const-str-array-decay.cl appears to be failing. The failing build can be seen in: https://lab.llvm.org/buildbot/#/builders/36/builds/33110

Thanks!
I will be able to address this issue within a few hours. If it can't wait, please revert it.

[barannikov88]

In D150520#4342429, @amyk wrote:

Hi, I just wanted to give a heads up that it appears that this patch is causing failures on some of the PPC buildbots.
In particular, the const-str-array-decay.cl appears to be failing. The failing build can be seen in: https://lab.llvm.org/buildbot/#/builders/36/builds/33110

Should be fixed by https://reviews.llvm.org/rGcc7dc90481d93734a56098470879189cb2e2c127

[amyk]

In D150520#4342943, @barannikov88 wrote:

In D150520#4342429, @amyk wrote:

Hi, I just wanted to give a heads up that it appears that this patch is causing failures on some of the PPC buildbots.
In particular, the const-str-array-decay.cl appears to be failing. The failing build can be seen in: https://lab.llvm.org/buildbot/#/builders/36/builds/33110

Should be fixed by https://reviews.llvm.org/rGcc7dc90481d93734a56098470879189cb2e2c127

Thanks for the quick fix!

clang/test/CodeGenCUDA/amdgpu-alias-undef-symbols.cu

 // REQUIRES: amdgpu-registered-target
 
-// RUN: %clang -Xclang -no-opaque-pointers -target x86_64-unknown-linux-gnu --offload-arch=gfx906 --cuda-device-only -nogpulib -nogpuinc -x hip -emit-llvm -S -o - %s \
+// RUN: %clang -target x86_64-unknown-linux-gnu --offload-arch=gfx906 --cuda-device-only -nogpulib -nogpuinc -x hip -emit-llvm -S -o - %s \
 // RUN:   -fgpu-rdc -O3 -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false | \
 // RUN:   FileCheck %s
 
 #include "Inputs/cuda.h"
 
-// CHECK: %struct.B = type { i8 }
 struct B {
 
-  // CHECK: @_ZN1BC1Ei = hidden unnamed_addr alias void (%struct.B*, i32), void (%struct.B*, i32)* @_ZN1BC2Ei
+  // CHECK: @_ZN1BC1Ei = hidden unnamed_addr alias void (ptr, i32), ptr @_ZN1BC2Ei
   __device__ B(int x);
 };
 
 __device__ B::B(int x) {
 }

clang/test/CodeGenCUDA/device-use-host-var.cu

-// RUN: %clang_cc1 -no-opaque-pointers -std=c++14 -triple amdgcn-amd-amdhsa \
+// RUN: %clang_cc1 -std=c++14 -triple amdgcn-amd-amdhsa \
 // RUN:   -fcuda-is-device -emit-llvm -o - -x hip %s | FileCheck %s
-// RUN: %clang_cc1 -no-opaque-pointers -std=c++14 -triple amdgcn-amd-amdhsa \
+// RUN: %clang_cc1 -std=c++14 -triple amdgcn-amd-amdhsa \
 // RUN:   -fcuda-is-device -emit-llvm -o - -x hip %s | FileCheck -check-prefix=NEG %s
 
 #include "Inputs/cuda.h"
 
 struct A {
   int x;
 };
 

 // NEG-NOT: @_Z1bIdE
 // NEG-NOT: @_ZL13var_host_only
 // NEG-NOT: {{^}}@{{.*}} = external
 
 // CHECK-LABEL: define{{.*}}@_Z7dev_funPiPPKi
 // CHECK: store i32 1
 // CHECK: store i32 2
 // CHECK: store i32 3
-// CHECK: load i8, i8* getelementptr {{.*}} @_ZL13constexpr_str.const
+// CHECK: load i8, ptr getelementptr {{.*}} @_ZL13constexpr_str.const
 // CHECK: store i32 4
 // CHECK: store i32 5
 // CHECK: store i32 6
-// CHECK: load i8, i8* getelementptr {{.*}} @_ZL9const_str
-// CHECK: store i32* {{.*}}@_ZL13constexpr_var
-// CHECK: store i32* getelementptr {{.*}} @_ZL16constexpr_struct
-// CHECK: store i32* getelementptr {{.*}} @_ZL15constexpr_array
-// CHECK: store i32* {{.*}}@_ZL9const_var
-// CHECK: store i32* getelementptr {{.*}} @_ZL12const_struct
-// CHECK: store i32* getelementptr {{.*}} @_ZL11const_array
+// CHECK: load i8, ptr getelementptr {{.*}} @_ZL9const_str
+// CHECK: store ptr {{.*}}@_ZL13constexpr_var
+// CHECK: store ptr {{.*}} @_ZL16constexpr_struct
+// CHECK: store ptr getelementptr {{.*}} @_ZL15constexpr_array
+// CHECK: store ptr {{.*}}@_ZL9const_var
+// CHECK: store ptr {{.*}} @_ZL12const_struct
+// CHECK: store ptr getelementptr {{.*}} @_ZL11const_array
 __device__ void dev_fun(int *out, const int **out2) {
   *out = constexpr_var;
   *out = constexpr_struct.x;
   *out = constexpr_array[3].x;
   *out = constexpr_str[3];
   *out = const_var;
   *out = const_struct.x;
   *out = const_array[3].x;

clang/test/CodeGenCUDA/managed-var.cu

-// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
 // RUN:   -emit-llvm -o - -x hip %s | FileCheck \
 // RUN:   -check-prefixes=COMMON,DEV,NORDC-D %s
 
-// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
 // RUN:   -emit-llvm -fgpu-rdc -cuid=abc -o - -x hip %s > %t.dev
 // RUN: cat %t.dev | FileCheck -check-prefixes=COMMON,DEV,RDC-D %s
 
-// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-gnu-linux -std=c++11 \
+// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \
 // RUN:   -emit-llvm -o - -x hip %s | FileCheck \
 // RUN:   -check-prefixes=COMMON,HOST,NORDC %s
 
-// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-gnu-linux -std=c++11 \
+// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \
 // RUN:   -emit-llvm -fgpu-rdc -cuid=abc -o - -x hip %s > %t.host
 // RUN: cat %t.host | FileCheck -check-prefixes=COMMON,HOST,RDC %s
 
 // Check device and host compilation use the same postfix for static
 // variable name.
 
 // RUN: cat %t.dev %t.host | FileCheck -check-prefix=POSTFIX %s
 
 #include "Inputs/cuda.h"
 
 struct vec {
   float x,y,z;
 };
 
 // DEV-DAG: @x.managed = addrspace(1) externally_initialized global i32 1, align 4
-// DEV-DAG: @x = addrspace(1) externally_initialized global i32 addrspace(1)* null
+// DEV-DAG: @x = addrspace(1) externally_initialized global ptr addrspace(1) null
 // NORDC-DAG: @x.managed = internal global i32 1
 // RDC-DAG: @x.managed = global i32 1
-// NORDC-DAG: @x = internal externally_initialized global i32* null
-// RDC-DAG: @x = externally_initialized global i32* null
+// NORDC-DAG: @x = internal externally_initialized global ptr null
+// RDC-DAG: @x = externally_initialized global ptr null
 // HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"x\00"
 __managed__ int x = 1;
 
 // DEV-DAG: @v.managed = addrspace(1) externally_initialized global [100 x %struct.vec] zeroinitializer, align 4
-// DEV-DAG: @v = addrspace(1) externally_initialized global [100 x %struct.vec] addrspace(1)* null
+// DEV-DAG: @v = addrspace(1) externally_initialized global ptr addrspace(1) null
 __managed__ vec v[100];
 
 // DEV-DAG: @v2.managed = addrspace(1) externally_initialized global <{ %struct.vec, [99 x %struct.vec] }> <{ %struct.vec { float 1.000000e+00, float 1.000000e+00, float 1.000000e+00 }, [99 x %struct.vec] zeroinitializer }>, align 4
-// DEV-DAG: @v2 = addrspace(1) externally_initialized global <{ %struct.vec, [99 x %struct.vec] }> addrspace(1)* null
+// DEV-DAG: @v2 = addrspace(1) externally_initialized global ptr addrspace(1) null
 __managed__ vec v2[100] = {{1, 1, 1}};
 
 // DEV-DAG: @ex.managed = external addrspace(1) global i32, align 4
-// DEV-DAG: @ex = external addrspace(1) externally_initialized global i32 addrspace(1)*
+// DEV-DAG: @ex = external addrspace(1) externally_initialized global ptr addrspace(1)
 // HOST-DAG: @ex.managed = external global i32
-// HOST-DAG: @ex = external externally_initialized global i32*
+// HOST-DAG: @ex = external externally_initialized global ptr
 extern __managed__ int ex;
 
 // NORDC-D-DAG: @_ZL2sx.managed = addrspace(1) externally_initialized global i32 1, align 4
-// NORDC-D-DAG: @_ZL2sx = addrspace(1) externally_initialized global i32 addrspace(1)* null
+// NORDC-D-DAG: @_ZL2sx = addrspace(1) externally_initialized global ptr addrspace(1) null
 // RDC-D-DAG: @_ZL2sx.static.[[HASH:.*]].managed = addrspace(1) externally_initialized global i32 1, align 4
-// RDC-D-DAG: @_ZL2sx.static.[[HASH]] = addrspace(1) externally_initialized global i32 addrspace(1)* null
+// RDC-D-DAG: @_ZL2sx.static.[[HASH]] = addrspace(1) externally_initialized global ptr addrspace(1) null
 // HOST-DAG: @_ZL2sx.managed = internal global i32 1
-// HOST-DAG: @_ZL2sx = internal externally_initialized global i32* null
+// HOST-DAG: @_ZL2sx = internal externally_initialized global ptr null
 // NORDC-DAG: @[[DEVNAMESX:[0-9]+]] = {{.*}}c"_ZL2sx\00"
 // RDC-DAG: @[[DEVNAMESX:[0-9]+]] = {{.*}}c"_ZL2sx.static.[[HASH:.*]]\00"
 
-// POSTFIX:  @_ZL2sx.static.[[HASH:.*]] = addrspace(1) externally_initialized global i32 addrspace(1)* null
+// POSTFIX:  @_ZL2sx.static.[[HASH:.*]] = addrspace(1) externally_initialized global ptr addrspace(1) null
 // POSTFIX: @[[DEVNAMESX:[0-9]+]] = {{.*}}c"_ZL2sx.static.[[HASH]]\00"
 static __managed__ int sx = 1;
 
 // DEV-DAG: @llvm.compiler.used
 // DEV-SAME-DAG: @x.managed
 // DEV-SAME-DAG: @x
 // DEV-SAME-DAG: @v.managed
 // DEV-SAME-DAG: @v
 // DEV-SAME-DAG: @_ZL2sx.managed
 // DEV-SAME-DAG: @_ZL2sx
 
 // Force ex and sx mitted in device compilation.
 __global__ void foo(int *z) {
   *z = x + ex + sx;
   v[1].x = 2;
 }
 
 // Force ex and sx emitted in host compilatioin.
 int foo2() {
   return ex + sx;
 }
 
 // COMMON-LABEL: define {{.*}}@_Z4loadv()
-// DEV:  %ld.managed = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(1)* @x, align 4
-// DEV:  %0 = addrspacecast i32 addrspace(1)* %ld.managed to i32*
-// DEV:  %1 = load i32, i32* %0, align 4
+// DEV:  %ld.managed = load ptr addrspace(1), ptr addrspace(1) @x, align 4
+// DEV:  %0 = addrspacecast ptr addrspace(1) %ld.managed to ptr
+// DEV:  %1 = load i32, ptr %0, align 4
 // DEV:  ret i32 %1
-// HOST:  %ld.managed = load i32*, i32** @x, align 4
-// HOST:  %0 = load i32, i32* %ld.managed, align 4
+// HOST:  %ld.managed = load ptr, ptr @x, align 4
+// HOST:  %0 = load i32, ptr %ld.managed, align 4
 // HOST:  ret i32 %0
 __device__ __host__ int load() {
   return x;
 }
 
 // COMMON-LABEL: define {{.*}}@_Z5storev()
-// DEV:  %ld.managed = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(1)* @x, align 4
-// DEV:  %0 = addrspacecast i32 addrspace(1)* %ld.managed to i32*
-// DEV:  store i32 2, i32* %0, align 4
-// HOST:  %ld.managed = load i32*, i32** @x, align 4
-// HOST:  store i32 2, i32* %ld.managed, align 4
+// DEV:  %ld.managed = load ptr addrspace(1), ptr addrspace(1) @x, align 4
+// DEV:  %0 = addrspacecast ptr addrspace(1) %ld.managed to ptr
+// DEV:  store i32 2, ptr %0, align 4
+// HOST:  %ld.managed = load ptr, ptr @x, align 4
+// HOST:  store i32 2, ptr %ld.managed, align 4
 __device__ __host__ void store() {
   x = 2;
 }
 
 // COMMON-LABEL: define {{.*}}@_Z10addr_takenv()
-// DEV:  %0 = addrspacecast i32 addrspace(1)* %ld.managed to i32*
-// DEV:  store i32* %0, i32** %p.ascast, align 8
-// DEV:  %1 = load i32*, i32** %p.ascast, align 8
-// DEV:  store i32 3, i32* %1, align 4
-// HOST:  %ld.managed = load i32*, i32** @x, align 4
-// HOST:  store i32* %ld.managed, i32** %p, align 8
-// HOST:  %0 = load i32*, i32** %p, align 8
-// HOST:  store i32 3, i32* %0, align 4
+// DEV:  %0 = addrspacecast ptr addrspace(1) %ld.managed to ptr
+// DEV:  store ptr %0, ptr %p.ascast, align 8
+// DEV:  %1 = load ptr, ptr %p.ascast, align 8
+// DEV:  store i32 3, ptr %1, align 4
+// HOST:  %ld.managed = load ptr, ptr @x, align 4
+// HOST:  store ptr %ld.managed, ptr %p, align 8
+// HOST:  %0 = load ptr, ptr %p, align 8
+// HOST:  store i32 3, ptr %0, align 4
 __device__ __host__ void addr_taken() {
   int *p = &x;
   *p = 3;
 }
 
 // HOST-LABEL: define {{.*}}@_Z5load2v()
-// HOST: %ld.managed = load [100 x %struct.vec]*, [100 x %struct.vec]** @v, align 16
-// HOST:  %0 = getelementptr inbounds [100 x %struct.vec], [100 x %struct.vec]* %ld.managed, i64 0, i64 1, i32 0
-// HOST:  %1 = load float, float* %0, align 4
+// HOST: %ld.managed = load ptr, ptr @v, align 16
+// HOST:  %0 = getelementptr inbounds [100 x %struct.vec], ptr %ld.managed, i64 0, i64 1
+// HOST:  %1 = load float, ptr %0, align 4
 // HOST:  ret float %1
 __device__ __host__ float load2() {
   return v[1].x;
 }
 
 // HOST-LABEL: define {{.*}}@_Z5load3v()
-// HOST:  %ld.managed = load <{ %struct.vec, [99 x %struct.vec] }>*, <{ %struct.vec, [99 x %struct.vec] }>** @v2, align 16
-// HOST:  %0 = bitcast <{ %struct.vec, [99 x %struct.vec] }>* %ld.managed to [100 x %struct.vec]*
-// HOST:  %1 = getelementptr inbounds [100 x %struct.vec], [100 x %struct.vec]* %0, i64 0, i64 1, i32 1
-// HOST:  %2 = load float, float* %1, align 4
-// HOST:  ret float %2
+// HOST:  %ld.managed = load ptr, ptr @v2, align 16
+// HOST:  %0 = getelementptr inbounds [100 x %struct.vec], ptr %ld.managed, i64 0, i64 1, i32 1
+// HOST:  %1 = load float, ptr %0, align 4
+// HOST:  ret float %1
 float load3() {
   return v2[1].y;
 }
 
 // HOST-LABEL: define {{.*}}@_Z11addr_taken2v()
-// HOST:  %ld.managed = load [100 x %struct.vec]*, [100 x %struct.vec]** @v, align 16
-// HOST:  %0 = getelementptr inbounds [100 x %struct.vec], [100 x %struct.vec]* %ld.managed, i64 0, i64 1, i32 0
-// HOST:  %1 = ptrtoint float* %0 to i64
-// HOST:  %ld.managed1 = load <{ %struct.vec, [99 x %struct.vec] }>*, <{ %struct.vec, [99 x %struct.vec] }>** @v2, align 16
-// HOST:  %2 = bitcast <{ %struct.vec, [99 x %struct.vec] }>* %ld.managed1 to [100 x %struct.vec]*
-// HOST:  %3 = getelementptr inbounds [100 x %struct.vec], [100 x %struct.vec]* %2, i64 0, i64 1, i32 1
-// HOST:  %4 = ptrtoint float* %3 to i64
-// HOST:  %5 = sub i64 %4, %1
-// HOST:  %sub.ptr.div = sdiv exact i64 %5, 4
+// HOST:  %ld.managed = load ptr, ptr @v, align 16
+// HOST:  %0 = getelementptr inbounds [100 x %struct.vec], ptr %ld.managed, i64 0, i64 1
+// HOST:  %1 = ptrtoint ptr %0 to i64
+// HOST:  %ld.managed1 = load ptr, ptr @v2, align 16
+// HOST:  %2 = getelementptr inbounds [100 x %struct.vec], ptr %ld.managed1, i64 0, i64 1, i32 1
+// HOST:  %3 = ptrtoint ptr %2 to i64
+// HOST:  %4 = sub i64 %3, %1
+// HOST:  %sub.ptr.div = sdiv exact i64 %4, 4
 // HOST:  %conv = sitofp i64 %sub.ptr.div to float
 // HOST:  ret float %conv
 float addr_taken2() {
   return (float)reinterpret_cast<long>(&(v2[1].y)-&(v[1].x));
 }
 
 // COMMON-LABEL: define {{.*}}@_Z5load4v()
-// DEV:  %ld.managed = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(1)* @ex, align 4
-// DEV:  %0 = addrspacecast i32 addrspace(1)* %ld.managed to i32*
-// DEV:  %1 = load i32, i32* %0, align 4
+// DEV:  %ld.managed = load ptr addrspace(1), ptr addrspace(1) @ex, align 4
+// DEV:  %0 = addrspacecast ptr addrspace(1) %ld.managed to ptr
+// DEV:  %1 = load i32, ptr %0, align 4
 // DEV:  ret i32 %1
-// HOST:  %ld.managed = load i32*, i32** @ex, align 4
-// HOST:  %0 = load i32, i32* %ld.managed, align 4
+// HOST:  %ld.managed = load ptr, ptr @ex, align 4
+// HOST:  %0 = load i32, ptr %ld.managed, align 4
 // HOST:  ret i32 %0
 __device__ __host__ int load4() {
   return ex;
 }
 
-// HOST-DAG: __hipRegisterManagedVar({{.*}}@x {{.*}}@x.managed {{.*}}@[[DEVNAMEX]]{{.*}}, i64 4, i32 4)
-// HOST-DAG: __hipRegisterManagedVar({{.*}}@_ZL2sx {{.*}}@_ZL2sx.managed {{.*}}@[[DEVNAMESX]]
-// HOST-NOT: __hipRegisterManagedVar({{.*}}@ex {{.*}}@ex.managed
-// HOST-DAG: declare void @__hipRegisterManagedVar(i8**, i8*, i8*, i8*, i64, i32)
+// HOST-DAG: __hipRegisterManagedVar({{.*}}, ptr @x, ptr @x.managed, ptr @[[DEVNAMEX]], i64 4, i32 4)
+// HOST-DAG: __hipRegisterManagedVar({{.*}}, ptr @_ZL2sx, ptr @_ZL2sx.managed, ptr @[[DEVNAMESX]]
+// HOST-NOT: __hipRegisterManagedVar({{.*}}, ptr @ex, ptr @ex.managed
+// HOST-DAG: declare void @__hipRegisterManagedVar(ptr, ptr, ptr, ptr, i64, i32)

clang/test/CodeGenCUDA/static-device-var-no-rdc.cu

 // REQUIRES: x86-registered-target
 // REQUIRES: amdgpu-registered-target
 
-// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
 // RUN:   -emit-llvm -o - -x hip %s | FileCheck -check-prefix=DEV %s
 
-// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-gnu-linux -std=c++11 \
+// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \
 // RUN:   -emit-llvm -o - -x hip %s | FileCheck -check-prefix=HOST %s
 
 // Negative tests.
 
-// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
 // RUN:   -emit-llvm -o - -x hip %s | FileCheck -check-prefix=DEV-NEG %s
 
-// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-gnu-linux -std=c++11 \
+// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \
 // RUN:   -emit-llvm -o - -x hip %s | FileCheck -check-prefix=HOST-NEG %s
 
 #include "Inputs/cuda.h"
 
 // Test function scope static device variable, which should not be externalized.
 // DEV-DAG: @_ZZ6kernelPiPPKiE1w = internal addrspace(4) constant i32 1
 // DEV-DAG: @_ZZ6kernelPiPPKiE21local_static_constant = internal addrspace(4) constant i32 42
 // DEV-DAG: @_ZZ6kernelPiPPKiE19local_static_device = internal addrspace(1) constant i32 43
void foo(const int **a) {
   getDeviceSymbol(&x5);
   getDeviceSymbol(&y);
   getDeviceSymbol(&w);
   z = 123;
   a[0] = &z2;
   decltype(u) tmp;
 }
 
-// HOST-DAG: __hipRegisterVar({{.*}}@_ZL1x {{.*}}@[[DEVNAMEX]]
-// HOST-DAG: __hipRegisterVar({{.*}}@_ZL1y {{.*}}@[[DEVNAMEY]]
-// HOST-DAG: __hipRegisterVar({{.*}}@_ZL1w {{.*}}@[[DEVNAMEW]]
+// HOST-DAG: __hipRegisterVar({{.*}}@_ZL1x, {{.*}}@[[DEVNAMEX]]
+// HOST-DAG: __hipRegisterVar({{.*}}@_ZL1y, {{.*}}@[[DEVNAMEY]]
+// HOST-DAG: __hipRegisterVar({{.*}}@_ZL1w, {{.*}}@[[DEVNAMEW]]
 // HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZL1u
 // HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZZ6kernelPiPPKiE1w
 // HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZZ6devfunPPKiE1p

clang/test/CodeGenCUDA/surface.cu

 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
 
-// RUN: %clang_cc1 -no-opaque-pointers -std=c++11 -fcuda-is-device -triple nvptx64-nvidia-cuda -emit-llvm -o - %s | FileCheck --check-prefix=DEVICE %s
+// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -triple nvptx64-nvidia-cuda -emit-llvm -o - %s | FileCheck --check-prefix=DEVICE %s
 // RUN: echo "GPU binary would be here" > %t
-// RUN: %clang_cc1 -no-opaque-pointers -std=c++11 -triple x86_64-unknown-linux-gnu -target-sdk-version=8.0 -fcuda-include-gpubinary %t -emit-llvm -o - %s | FileCheck --check-prefix=HOST %s
+// RUN: %clang_cc1 -std=c++11 -triple x86_64-unknown-linux-gnu -target-sdk-version=8.0 -fcuda-include-gpubinary %t -emit-llvm -o - %s | FileCheck --check-prefix=HOST %s
 
 struct surfaceReference {
   int desc;
 };
 
 template <typename T, int dim = 1>
 struct __attribute__((device_builtin_surface_type)) surface : public surfaceReference {
 };
 
 // Partial specialization over `void`.
 template<int dim>
 struct __attribute__((device_builtin_surface_type)) surface<void, dim> : public surfaceReference {
 };
 
 // On the device side, surface references are represented as `i64` handles.
 // DEVICE: @surf ={{.*}} addrspace(1) externally_initialized global i64 undef, align 4
 // On the host side, they remain in the original type.
 // HOST: @surf = internal global %struct.surface
 // HOST: @0 = private unnamed_addr constant [5 x i8] c"surf\00"
 surface<void, 2> surf;
 
 __attribute__((device)) int suld_2d_zero(surface<void, 2>, int, int) asm("llvm.nvvm.suld.2d.i32.zero");
 
 // DEVICE-LABEL: i32 @_Z3fooii(i32 noundef %x, i32 noundef %y)
-// DEVICE: call i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)* @surf)
+// DEVICE: call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @surf)
 // DEVICE: call noundef i32 @llvm.nvvm.suld.2d.i32.zero(i64 %{{.*}}, i32 noundef %{{.*}}, i32 noundef %{{.*}})
 __attribute__((device)) int foo(int x, int y) {
   return suld_2d_zero(surf, x, y);
 }
 
 // HOST: define internal void @[[PREFIX:__cuda]]_register_globals
 // Texture references need registering with correct arguments.
-// HOST: call void @[[PREFIX]]RegisterSurface(i8** %0, i8*{{.*}}({{.*}}@surf{{.*}}), i8*{{.*}}({{.*}}@0{{.*}}), i8*{{.*}}({{.*}}@0{{.*}}), i32 2, i32 0)
+// HOST: call void @[[PREFIX]]RegisterSurface(ptr %0, ptr @surf, ptr @0, ptr @0, i32 2, i32 0)
 
 // They also need annotating in metadata.
-// DEVICE: !0 = !{i64 addrspace(1)* @surf, !"surface", i32 1}
+// DEVICE: !0 = !{ptr addrspace(1) @surf, !"surface", i32 1}

clang/test/CodeGenCUDA/usual-deallocators.cu

-// RUN: %clang_cc1 -no-opaque-pointers %s --std=c++11 -triple nvptx-unknown-unknown -fcuda-is-device \
+// RUN: %clang_cc1 %s --std=c++11 -triple nvptx-unknown-unknown -fcuda-is-device \
 // RUN:   -emit-llvm -o - | FileCheck %s --check-prefixes=COMMON,DEVICE
-// RUN: %clang_cc1 -no-opaque-pointers %s --std=c++11 -triple nvptx-unknown-unknown \
+// RUN: %clang_cc1 %s --std=c++11 -triple nvptx-unknown-unknown \
 // RUN:   -emit-llvm -o - | FileCheck %s --check-prefixes=COMMON,HOST
-// RUN: %clang_cc1 -no-opaque-pointers %s --std=c++17 -triple nvptx-unknown-unknown -fcuda-is-device \
+// RUN: %clang_cc1 %s --std=c++17 -triple nvptx-unknown-unknown -fcuda-is-device \
 // RUN:   -emit-llvm -o - | FileCheck %s --check-prefixes=COMMON,DEVICE
-// RUN: %clang_cc1 -no-opaque-pointers %s --std=c++17 -triple nvptx-unknown-unknown \
+// RUN: %clang_cc1 %s --std=c++17 -triple nvptx-unknown-unknown \
 // RUN:   -emit-llvm -o - | FileCheck %s --check-prefixes=COMMON,HOST
 
 #include "Inputs/cuda.h"
 extern "C" __host__ void host_fn();
 extern "C" __device__ void dev_fn();
 extern "C" __host__ __device__ void hd_fn();
 
 // Destructors are handled a bit differently, compared to regular functions.

 }
 
 // Make sure we call the right variant of usual deallocator.
 __host__ __device__ void tests_hd(void *t) {
   // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H1D1EvPv
   // COMMON: call void @_ZN4H1D1dlEPv
   test_hd<H1D1>(t);
   // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H1D2EvPv
-  // DEVICE: call void @_ZN4H1D2dlEPvj(i8* noundef {{.*}}, i32 noundef 1)
-  // HOST:   call void @_ZN4H1D2dlEPv(i8* noundef {{.*}})
+  // DEVICE: call void @_ZN4H1D2dlEPvj(ptr noundef {{.*}}, i32 noundef 1)
+  // HOST:   call void @_ZN4H1D2dlEPv(ptr noundef {{.*}})
   test_hd<H1D2>(t);
   // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H2D1EvPv
-  // DEVICE: call void @_ZN4H2D1dlEPv(i8* {{.*}})
-  // HOST:   call void @_ZN4H2D1dlEPvj(i8* noundef %3, i32 noundef 1)
+  // DEVICE: call void @_ZN4H2D1dlEPv(ptr {{.*}})
+  // HOST:   call void @_ZN4H2D1dlEPvj(ptr noundef {{.*}}, i32 noundef 1)
   test_hd<H2D1>(t);
   // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H2D2EvPv
-  // COMMON: call void @_ZN4H2D2dlEPvj(i8* noundef {{.*}}, i32 noundef 1)
+  // COMMON: call void @_ZN4H2D2dlEPvj(ptr noundef {{.*}}, i32 noundef 1)
   test_hd<H2D2>(t);
   // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1D1D2EvPv
-  // COMMON: call void @_ZN6H1D1D2dlEPv(i8* noundef %3)
+  // COMMON: call void @_ZN6H1D1D2dlEPv(ptr noundef {{.*}})
   test_hd<H1D1D2>(t);
   // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1H2D1EvPv
-  // COMMON: call void @_ZN6H1H2D1dlEPv(i8* {{.*}})
+  // COMMON: call void @_ZN6H1H2D1dlEPv(ptr {{.*}})
   test_hd<H1H2D1>(t);
   // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1H2D2EvPv
-  // DEVICE: call void @_ZN6H1H2D2dlEPvj(i8* noundef {{.*}}, i32 noundef 1)
-  // HOST:   call void @_ZN6H1H2D2dlEPv(i8* noundef {{.*}})
+  // DEVICE: call void @_ZN6H1H2D2dlEPvj(ptr noundef {{.*}}, i32 noundef 1)
+  // HOST:   call void @_ZN6H1H2D2dlEPv(ptr noundef {{.*}})
   test_hd<H1H2D2>(t);
   // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI8H1H2D1D2EvPv
-  // COMMON: call void @_ZN8H1H2D1D2dlEPv(i8* {{.*}})
+  // COMMON: call void @_ZN8H1H2D1D2dlEPv(ptr {{.*}})
   test_hd<H1H2D1D2>(t);
 }
 
 // Make sure that we've generated the kernel used by A::~A.
 // DEVICE-LABEL: define void @_Z1fIiEvT_
 
 // Make sure we've picked deallocator for the correct side of compilation.
 
-// COMMON-LABEL: define  linkonce_odr void @_ZN4H1D1dlEPv(i8* noundef %0)
+// COMMON-LABEL: define  linkonce_odr void @_ZN4H1D1dlEPv(ptr noundef %0)
 // DEVICE: call void @dev_fn()
 // HOST:   call void @host_fn()
 
-// DEVICE-LABEL: define  linkonce_odr void @_ZN4H1D2dlEPvj(i8* noundef %0, i32 noundef %1)
+// DEVICE-LABEL: define  linkonce_odr void @_ZN4H1D2dlEPvj(ptr noundef %0, i32 noundef %1)
 // DEVICE: call void @dev_fn()
-// HOST-LABEL: define linkonce_odr void @_ZN4H1D2dlEPv(i8* noundef %0)
+// HOST-LABEL: define linkonce_odr void @_ZN4H1D2dlEPv(ptr noundef %0)
 // HOST: call void @host_fn()
 
-// DEVICE-LABEL: define  linkonce_odr void @_ZN4H2D1dlEPv(i8* noundef %0)
+// DEVICE-LABEL: define  linkonce_odr void @_ZN4H2D1dlEPv(ptr noundef %0)
 // DEVICE: call void @dev_fn()
-// HOST-LABEL:  define linkonce_odr void @_ZN4H2D1dlEPvj(i8* noundef %0, i32 noundef %1)
+// HOST-LABEL:  define linkonce_odr void @_ZN4H2D1dlEPvj(ptr noundef %0, i32 noundef %1)
 // HOST: call void @host_fn()
 
-// COMMON-LABEL: define  linkonce_odr void @_ZN4H2D2dlEPvj(i8* noundef %0, i32 noundef %1)
+// COMMON-LABEL: define  linkonce_odr void @_ZN4H2D2dlEPvj(ptr noundef %0, i32 noundef %1)
 // DEVICE: call void @dev_fn()
 // HOST: call void @host_fn()
 
-// COMMON-LABEL: define  linkonce_odr void @_ZN6H1D1D2dlEPv(i8* noundef %0)
+// COMMON-LABEL: define  linkonce_odr void @_ZN6H1D1D2dlEPv(ptr noundef %0)
 // DEVICE: call void @dev_fn()
 // HOST: call void @host_fn()
 
-// COMMON-LABEL: define  linkonce_odr void @_ZN6H1H2D1dlEPv(i8* noundef %0)
+// COMMON-LABEL: define  linkonce_odr void @_ZN6H1H2D1dlEPv(ptr noundef %0)
 // DEVICE: call void @dev_fn()
 // HOST: call void @host_fn()
 
-// DEVICE-LABEL: define  linkonce_odr void @_ZN6H1H2D2dlEPvj(i8* noundef %0, i32 noundef %1)
+// DEVICE-LABEL: define  linkonce_odr void @_ZN6H1H2D2dlEPvj(ptr noundef %0, i32 noundef %1)
 // DEVICE: call void @dev_fn()
-// HOST-LABEL: define linkonce_odr void @_ZN6H1H2D2dlEPv(i8* noundef %0)
+// HOST-LABEL: define linkonce_odr void @_ZN6H1H2D2dlEPv(ptr noundef %0)
 // HOST: call void @host_fn()
 
-// COMMON-LABEL: define  linkonce_odr void @_ZN8H1H2D1D2dlEPv(i8* noundef %0)
+// COMMON-LABEL: define  linkonce_odr void @_ZN8H1H2D1D2dlEPv(ptr noundef %0)
 // DEVICE: call void @dev_fn()
 // HOST: call void @host_fn()
 
-// DEVICE: !0 = !{void (i32)* @_Z1fIiEvT_, !"kernel", i32 1}
+// DEVICE: !0 = !{ptr @_Z1fIiEvT_, !"kernel", i32 1}

clang/test/CodeGenCUDA/vtbl.cu

-// RUN: %clang_cc1 -no-opaque-pointers -fcuda-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx906 \
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx906 \
 // RUN:   -emit-llvm -o - %s | FileCheck %s
 
 #include "Inputs/cuda.h"
 
-// CHECK-LABEL: define {{.*}}@_ZN1AC2Ev(%struct.A* noundef nonnull align 8 dereferenceable(8) %this)
-// CHECK: store %struct.A* %this, %struct.A** %this.addr.ascast
-// CHECK: %this1 = load %struct.A*, %struct.A** %this.addr.ascast
-// CHECK: %[[VTFIELD:.*]] = bitcast %struct.A* %this1 to i32 (...)* addrspace(1)**
-// CHECK: store i32 (...)* addrspace(1)* bitcast{{.*}} @_ZTV1A{{.*}}, i32 (...)* addrspace(1)** %[[VTFIELD]]
+// CHECK-LABEL: define {{.*}}@_ZN1AC2Ev(ptr noundef nonnull align 8 dereferenceable(8) %this)
+// CHECK: store ptr %this, ptr %this.addr.ascast
+// CHECK: %this1 = load ptr, ptr %this.addr.ascast
+// CHECK: store ptr addrspace(1) {{.*}} @_ZTV1A{{.*}}, ptr %this1
 struct A {
   __device__ virtual void vf() {}
 };
 
 __global__ void kern() {
   A a;
 }

clang/test/CodeGenCoroutines/coro-params.cpp

 // Verifies that parameters are copied with move constructors
 // Verifies that parameter copies are destroyed
 // Vefifies that parameter copies are used in the body of the coroutine
 // Verifies that parameter copies are used to construct the promise type, if that type has a matching constructor
-// RUN: %clang_cc1 -no-opaque-pointers -std=c++20 -triple=x86_64-unknown-linux-gnu -emit-llvm -o - %s -disable-llvm-passes -fexceptions | FileCheck %s
+// RUN: %clang_cc1 -std=c++20 -triple=x86_64-unknown-linux-gnu -emit-llvm -o - %s -disable-llvm-passes -fexceptions | FileCheck %s
 
 namespace std {
 template <typename... T> struct coroutine_traits;
 
 template <class Promise = void> struct coroutine_handle {
   coroutine_handle() = default;
   static coroutine_handle from_address(void *) noexcept;
 };
struct MoveAndCopy {
   MoveAndCopy(const MoveAndCopy&)noexcept;
   MoveAndCopy(MoveAndCopy&&) noexcept;
   ~MoveAndCopy();
 };
 
 void consume(int,int,int) noexcept;
 
 // TODO: Add support for CopyOnly params
-// CHECK: define{{.*}} void @_Z1fi8MoveOnly11MoveAndCopy(i32 noundef %val, %struct.MoveOnly* noundef %[[MoParam:.+]], %struct.MoveAndCopy* noundef %[[McParam:.+]]) #0 personality i8* bitcast (i32 (...)* @__gxx_personality_v0 to i8*
+// CHECK: define{{.*}} void @_Z1fi8MoveOnly11MoveAndCopy(i32 noundef %val, ptr noundef %[[MoParam:.+]], ptr noundef %[[McParam:.+]]) #0 personality ptr @__gxx_personality_v0
 void f(int val, MoveOnly moParam, MoveAndCopy mcParam) {
   // CHECK: %[[MoCopy:.+]] = alloca %struct.MoveOnly,
   // CHECK: %[[McCopy:.+]] = alloca %struct.MoveAndCopy,
-  // CHECK: store i32 %val, i32* %[[ValAddr:.+]]
+  // CHECK: store i32 %val, ptr %[[ValAddr:.+]]
 
-  // CHECK: call i8* @llvm.coro.begin(
-  // CHECK: call void @_ZN8MoveOnlyC1EOS_(%struct.MoveOnly* {{[^,]*}} %[[MoCopy]], %struct.MoveOnly* noundef nonnull align 4 dereferenceable(4) %[[MoParam]])
-  // CHECK-NEXT: bitcast %struct.MoveAndCopy* %[[McCopy]] to i8*
-  // CHECK-NEXT: call void @llvm.lifetime.start.p0i8(
-  // CHECK-NEXT: call void @_ZN11MoveAndCopyC1EOS_(%struct.MoveAndCopy* {{[^,]*}} %[[McCopy]], %struct.MoveAndCopy* noundef nonnull align 4 dereferenceable(4) %[[McParam]]) #
-  // CHECK-NEXT: bitcast %"struct.std::coroutine_traits<void, int, MoveOnly, MoveAndCopy>::promise_type"* %__promise to i8*
-  // CHECK-NEXT: call void @llvm.lifetime.start.p0i8(
+  // CHECK: call ptr @llvm.coro.begin(
+  // CHECK: call void @_ZN8MoveOnlyC1EOS_(ptr {{[^,]*}} %[[MoCopy]], ptr noundef nonnull align 4 dereferenceable(4) %[[MoParam]])
+  // CHECK-NEXT: call void @llvm.lifetime.start.p0(
+  // CHECK-NEXT: call void @_ZN11MoveAndCopyC1EOS_(ptr {{[^,]*}} %[[McCopy]], ptr noundef nonnull align 4 dereferenceable(4) %[[McParam]]) #
+  // CHECK-NEXT: call void @llvm.lifetime.start.p0(
   // CHECK-NEXT: invoke void @_ZNSt16coroutine_traitsIJvi8MoveOnly11MoveAndCopyEE12promise_typeC1Ev(
 
   // CHECK: call void @_ZN14suspend_always12await_resumeEv(
-  // CHECK: %[[IntParam:.+]] = load i32, i32* %{{.*}}
-  // CHECK: %[[MoGep:.+]] = getelementptr inbounds %struct.MoveOnly, %struct.MoveOnly* %[[MoCopy]], i32 0, i32 0
-  // CHECK: %[[MoVal:.+]] = load i32, i32* %[[MoGep]]
-  // CHECK: %[[McGep:.+]] =  getelementptr inbounds %struct.MoveAndCopy, %struct.MoveAndCopy* %[[McCopy]], i32 0, i32 0
-  // CHECK: %[[McVal:.+]] = load i32, i32* %[[McGep]]
+  // CHECK: %[[IntParam:.+]] = load i32, ptr %{{.*}}
+  // CHECK: %[[MoGep:.+]] = getelementptr inbounds %struct.MoveOnly, ptr %[[MoCopy]], i32 0, i32 0
+  // CHECK: %[[MoVal:.+]] = load i32, ptr %[[MoGep]]
+  // CHECK: %[[McGep:.+]] =  getelementptr inbounds %struct.MoveAndCopy, ptr %[[McCopy]], i32 0, i32 0
+  // CHECK: %[[McVal:.+]] = load i32, ptr %[[McGep]]
   // CHECK: call void @_Z7consumeiii(i32 noundef %[[IntParam]], i32 noundef %[[MoVal]], i32 noundef %[[McVal]])
 
   consume(val, moParam.val, mcParam.val);
   co_return;
 
   // Skip to final suspend:
   // CHECK: call void @_ZNSt16coroutine_traitsIJvi8MoveOnly11MoveAndCopyEE12promise_type13final_suspendEv(
   // CHECK: call void @_ZN14suspend_always12await_resumeEv(
 
   // Destroy promise, then parameter copies:
-  // CHECK: call void @_ZNSt16coroutine_traitsIJvi8MoveOnly11MoveAndCopyEE12promise_typeD1Ev(%"struct.std::coroutine_traits<void, int, MoveOnly, MoveAndCopy>::promise_type"* {{[^,]*}} %__promise)
-  // CHECK-NEXT: bitcast %"struct.std::coroutine_traits<void, int, MoveOnly, MoveAndCopy>::promise_type"* %__promise to i8*
-  // CHECK-NEXT: call void @llvm.lifetime.end.p0i8(
-  // CHECK-NEXT: call void @_ZN11MoveAndCopyD1Ev(%struct.MoveAndCopy* {{[^,]*}} %[[McCopy]])
-  // CHECK-NEXT: bitcast %struct.MoveAndCopy* %[[McCopy]] to i8*
-  // CHECK-NEXT: call void @llvm.lifetime.end.p0i8(
-  // CHECK-NEXT: call void @_ZN8MoveOnlyD1Ev(%struct.MoveOnly* {{[^,]*}} %[[MoCopy]]
-  // CHECK-NEXT: bitcast %struct.MoveOnly* %[[MoCopy]] to i8*
-  // CHECK-NEXT: call void @llvm.lifetime.end.p0i8(
-  // CHECK-NEXT: bitcast i32* %{{.+}} to i8*
-  // CHECK-NEXT: call void @llvm.lifetime.end.p0i8(
-  // CHECK-NEXT: call i8* @llvm.coro.free(
+  // CHECK: call void @_ZNSt16coroutine_traitsIJvi8MoveOnly11MoveAndCopyEE12promise_typeD1Ev(ptr {{[^,]*}} %__promise)
+  // CHECK-NEXT: call void @llvm.lifetime.end.p0(
+  // CHECK-NEXT: call void @_ZN11MoveAndCopyD1Ev(ptr {{[^,]*}} %[[McCopy]])
+  // CHECK-NEXT: call void @llvm.lifetime.end.p0(
+  // CHECK-NEXT: call void @_ZN8MoveOnlyD1Ev(ptr {{[^,]*}} %[[MoCopy]]
+  // CHECK-NEXT: call void @llvm.lifetime.end.p0(
+  // CHECK-NEXT: call void @llvm.lifetime.end.p0(
+  // CHECK-NEXT: call ptr @llvm.coro.free(
 }
 
-// CHECK-LABEL: void @_Z16dependent_paramsI1A1BEvT_T0_S3_(%struct.A* noundef %x, %struct.B* noundef %0, %struct.B* noundef %y)
+// CHECK-LABEL: void @_Z16dependent_paramsI1A1BEvT_T0_S3_(ptr noundef %x, ptr noundef %0, ptr noundef %y)
 template <typename T, typename U>
 void dependent_params(T x, U, U y) {
   // CHECK: %[[x_copy:.+]] = alloca %struct.A,
   // CHECK-NEXT: %[[unnamed_copy:.+]] = alloca %struct.B
   // CHECK-NEXT: %[[y_copy:.+]] = alloca %struct.B
 
-  // CHECK: call i8* @llvm.coro.begin
-  // CHECK-NEXT: bitcast %struct.A* %[[x_copy]] to i8*
-  // CHECK-NEXT: call void @llvm.lifetime.start.p0i8(
-  // CHECK-NEXT: call void @_ZN1AC1EOS_(%struct.A* {{[^,]*}} %[[x_copy]], %struct.A* noundef nonnull align 4 dereferenceable(512) %x)
-  // CHECK-NEXT: bitcast %struct.B* %[[unnamed_copy]] to i8*
-  // CHECK-NEXT: call void @llvm.lifetime.start.p0i8(
-  // CHECK-NEXT: call void @_ZN1BC1EOS_(%struct.B* {{[^,]*}} %[[unnamed_copy]], %struct.B* noundef nonnull align 4 dereferenceable(512) %0)
-  // CHECK-NEXT: bitcast %struct.B* %[[y_copy]] to i8*
-  // CHECK-NEXT: call void @llvm.lifetime.start.p0i8(
-  // CHECK-NEXT: call void @_ZN1BC1EOS_(%struct.B* {{[^,]*}} %[[y_copy]], %struct.B* noundef nonnull align 4 dereferenceable(512) %y)
-  // CHECK-NEXT: bitcast %"struct.std::coroutine_traits<void, A, B, B>::promise_type"* %__promise to i8*
-  // CHECK-NEXT: call void @llvm.lifetime.start.p0i8(
+  // CHECK: call ptr @llvm.coro.begin
+  // CHECK-NEXT: call void @llvm.lifetime.start.p0(
+  // CHECK-NEXT: call void @_ZN1AC1EOS_(ptr {{[^,]*}} %[[x_copy]], ptr noundef nonnull align 4 dereferenceable(512) %x)
+  // CHECK-NEXT: call void @llvm.lifetime.start.p0(
+  // CHECK-NEXT: call void @_ZN1BC1EOS_(ptr {{[^,]*}} %[[unnamed_copy]], ptr noundef nonnull align 4 dereferenceable(512) %0)
+  // CHECK-NEXT: call void @llvm.lifetime.start.p0(
+  // CHECK-NEXT: call void @_ZN1BC1EOS_(ptr {{[^,]*}} %[[y_copy]], ptr noundef nonnull align 4 dereferenceable(512) %y)
+  // CHECK-NEXT: call void @llvm.lifetime.start.p0(
   // CHECK-NEXT: invoke void @_ZNSt16coroutine_traitsIJv1A1BS1_EE12promise_typeC1Ev(
 
   co_return;
 }
 
 struct A {
   int WontFitIntoRegisterForSure[128];
   A();
struct promise_type {
     suspend_always final_suspend() noexcept { return {}; }
     void return_void() {}
     void unhandled_exception() {}
   };
 };
 
 // CHECK-LABEL: void @_Z38coroutine_matching_promise_constructor28promise_matching_constructorifd(i32 noundef %0, float noundef %1, double noundef %2)
 void coroutine_matching_promise_constructor(promise_matching_constructor, int, float, double) {
-  // CHECK: %[[INT:.+]] = load i32, i32* %5, align 4
-  // CHECK: %[[FLOAT:.+]] = load float, float* %6, align 4
-  // CHECK: %[[DOUBLE:.+]] = load double, double* %7, align 8
-  // CHECK: invoke void @_ZNSt16coroutine_traitsIJv28promise_matching_constructorifdEE12promise_typeC1ES0_ifd(%"struct.std::coroutine_traits<void, promise_matching_constructor, int, float, double>::promise_type"* {{[^,]*}} %__promise, i32 noundef %[[INT]], float noundef %[[FLOAT]], double noundef %[[DOUBLE]])
+  // CHECK: %[[INT:.+]] = load i32, ptr %5, align 4
+  // CHECK: %[[FLOAT:.+]] = load float, ptr %6, align 4
+  // CHECK: %[[DOUBLE:.+]] = load double, ptr %7, align 8
+  // CHECK: invoke void @_ZNSt16coroutine_traitsIJv28promise_matching_constructorifdEE12promise_typeC1ES0_ifd(ptr {{[^,]*}} %__promise, i32 noundef %[[INT]], float noundef %[[FLOAT]], double noundef %[[DOUBLE]])
   co_return;
 }
 
 struct some_class;
 
 struct method {};
 
 template <typename... Args> struct std::coroutine_traits<method, Args...> {
   struct promise_type {
     promise_type(some_class&, float);
     method get_return_object();
     suspend_always initial_suspend();
     suspend_always final_suspend() noexcept;
     void return_void();
     void unhandled_exception();
   };
 };
 
 struct some_class {
   method good_coroutine_calls_custom_constructor(float);
 };
 
-// CHECK-LABEL: define{{.*}} void @_ZN10some_class39good_coroutine_calls_custom_constructorEf(%struct.some_class*
+// CHECK-LABEL: define{{.*}} void @_ZN10some_class39good_coroutine_calls_custom_constructorEf(ptr
 method some_class::good_coroutine_calls_custom_constructor(float) {
-  // CHECK: invoke void @_ZNSt16coroutine_traitsIJ6methodR10some_classfEE12promise_typeC1ES2_f(%"struct.std::coroutine_traits<method, some_class &, float>::promise_type"* {{[^,]*}} %__promise, %struct.some_class* noundef nonnull align 1 dereferenceable(1) %{{.+}}, float
+  // CHECK: invoke void @_ZNSt16coroutine_traitsIJ6methodR10some_classfEE12promise_typeC1ES2_f(ptr {{[^,]*}} %__promise, ptr noundef nonnull align 1 dereferenceable(1) %{{.+}}, float
   co_return;
 }

clang/test/CodeGenHIP/maybe_undef-attr-verify.hip

-// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -target-cpu gfx906 -x hip -fcuda-is-device -emit-llvm  %s \
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx906 -x hip -fcuda-is-device -emit-llvm  %s \
 // RUN:   -o - | FileCheck %s
 
 // CHECK: define dso_local amdgpu_kernel void @_Z13shufflekernelv()
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP1:%.*]] = alloca i32, align 4, addrspace(5)
 // CHECK-NEXT:    [[TMP2:%.*]] = alloca i32, align 4, addrspace(5)
-// CHECK-NEXT:    [[TMP3:%.*]] = addrspacecast i32 addrspace(5)* [[TMP1:%.*]] to i32*
-// CHECK-NEXT:    [[TMP4:%.*]] = addrspacecast i32 addrspace(5)* [[TMP2:%.*]] to i32*
-// CHECK-NEXT:    [[TMP5:%.*]] = load i32, i32* [[TMP3:%.*]], align 4
+// CHECK-NEXT:    [[TMP3:%.*]] = addrspacecast ptr addrspace(5) [[TMP1:%.*]] to ptr
+// CHECK-NEXT:    [[TMP4:%.*]] = addrspacecast ptr addrspace(5) [[TMP2:%.*]] to ptr
+// CHECK-NEXT:    [[TMP5:%.*]] = load i32, ptr [[TMP3:%.*]], align 4
 // CHECK-NEXT:    [[TMP6:%.*]] = freeze i32 [[TMP5:%.*]]
 // CHECK-NEXT:    %call = call noundef i32 @_Z11__shfl_synciii(i32 noundef [[TMP6:%.*]], i32 noundef 64, i32 noundef 0) #4
-// CHECK-NEXT:    store i32 %call, i32* [[TMP4:%.*]], align 4
+// CHECK-NEXT:    store i32 %call, ptr [[TMP4:%.*]], align 4
 // CHECK-NEXT:  ret void
 
 // CHECK: define linkonce_odr noundef i32 @_Z11__shfl_synciii(i32 noundef [[TMP1:%.*]], i32 noundef [[TMP2:%.*]], i32 noundef [[TMP3:%.*]])
 
 #define __global__ __attribute__((global))
 #define __device__ __attribute__((device))
 #define __maybe_undef __attribute__((maybe_undef))
 #define WARP_SIZE 64

clang/test/CodeGenOpenCL/amdgpu-alignment.cl

 // REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-unknown-unknown -S -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
-// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-unknown-unknown-opencl -S -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -S -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown-opencl -S -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
 
 #pragma OPENCL EXTENSION cl_khr_fp64 : enable
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
 typedef char __attribute__((ext_vector_type(2))) char2;
 typedef char __attribute__((ext_vector_type(3))) char3;
 typedef char __attribute__((ext_vector_type(4))) char4;
 typedef char __attribute__((ext_vector_type(8))) char8;

 // CHECK: @local_memory_alignment_global.lds_v2f64 = internal addrspace(3) global [4 x <2 x double>] undef, align 16
 // CHECK: @local_memory_alignment_global.lds_v3f64 = internal addrspace(3) global [4 x <3 x double>] undef, align 32
 // CHECK: @local_memory_alignment_global.lds_v4f64 = internal addrspace(3) global [4 x <4 x double>] undef, align 32
 // CHECK: @local_memory_alignment_global.lds_v8f64 = internal addrspace(3) global [4 x <8 x double>] undef, align 64
 // CHECK: @local_memory_alignment_global.lds_v16f64 = internal addrspace(3) global [4 x <16 x double>] undef, align 128
 
 
 // CHECK-LABEL: @local_memory_alignment_global(
-// CHECK: store volatile i8 0, i8 addrspace(3)* getelementptr inbounds ([4 x i8], [4 x i8] addrspace(3)* @local_memory_alignment_global.lds_i8, i64 0, i64 0), align 1
-// CHECK: store volatile <2 x i8> zeroinitializer, <2 x i8> addrspace(3)* getelementptr inbounds ([4 x <2 x i8>], [4 x <2 x i8>] addrspace(3)* @local_memory_alignment_global.lds_v2i8, i64 0, i64 0), align 2
-// CHECK: store volatile <4 x i8> <i8 0, i8 0, i8 0, i8 undef>, <4 x i8> addrspace(3)* bitcast ([4 x <3 x i8>] addrspace(3)* @local_memory_alignment_global.lds_v3i8 to <4 x i8> addrspace(3)*), align 4
-// CHECK: store volatile <4 x i8> zeroinitializer, <4 x i8> addrspace(3)* getelementptr inbounds ([4 x <4 x i8>], [4 x <4 x i8>] addrspace(3)* @local_memory_alignment_global.lds_v4i8, i64 0, i64 0), align 4
-// CHECK: store volatile <8 x i8> zeroinitializer, <8 x i8> addrspace(3)* getelementptr inbounds ([4 x <8 x i8>], [4 x <8 x i8>] addrspace(3)* @local_memory_alignment_global.lds_v8i8, i64 0, i64 0), align 8
-// CHECK: store volatile <16 x i8> zeroinitializer, <16 x i8> addrspace(3)* getelementptr inbounds ([4 x <16 x i8>], [4 x <16 x i8>] addrspace(3)* @local_memory_alignment_global.lds_v16i8, i64 0, i64 0), align 16
-// CHECK: store volatile i16 0, i16 addrspace(3)* getelementptr inbounds ([4 x i16], [4 x i16] addrspace(3)* @local_memory_alignment_global.lds_i16, i64 0, i64 0), align 2
-// CHECK: store volatile <2 x i16> zeroinitializer, <2 x i16> addrspace(3)* getelementptr inbounds ([4 x <2 x i16>], [4 x <2 x i16>] addrspace(3)* @local_memory_alignment_global.lds_v2i16, i64 0, i64 0), align 4
-// CHECK: store volatile <4 x i16> <i16 0, i16 0, i16 0, i16 undef>, <4 x i16> addrspace(3)* bitcast ([4 x <3 x i16>] addrspace(3)* @local_memory_alignment_global.lds_v3i16 to <4 x i16> addrspace(3)*), align 8
-// CHECK: store volatile <4 x i16> zeroinitializer, <4 x i16> addrspace(3)* getelementptr inbounds ([4 x <4 x i16>], [4 x <4 x i16>] addrspace(3)* @local_memory_alignment_global.lds_v4i16, i64 0, i64 0), align 8
-// CHECK: store volatile <8 x i16> zeroinitializer, <8 x i16> addrspace(3)* getelementptr inbounds ([4 x <8 x i16>], [4 x <8 x i16>] addrspace(3)* @local_memory_alignment_global.lds_v8i16, i64 0, i64 0), align 16
-// CHECK: store volatile <16 x i16> zeroinitializer, <16 x i16> addrspace(3)* getelementptr inbounds ([4 x <16 x i16>], [4 x <16 x i16>] addrspace(3)* @local_memory_alignment_global.lds_v16i16, i64 0, i64 0), align 32
-// CHECK: store volatile i32 0, i32 addrspace(3)* getelementptr inbounds ([4 x i32], [4 x i32] addrspace(3)* @local_memory_alignment_global.lds_i32, i64 0, i64 0), align 4
-// CHECK: store volatile <2 x i32> zeroinitializer, <2 x i32> addrspace(3)* getelementptr inbounds ([4 x <2 x i32>], [4 x <2 x i32>] addrspace(3)* @local_memory_alignment_global.lds_v2i32, i64 0, i64 0), align 8
-// CHECK: store volatile <4 x i32> <i32 0, i32 0, i32 0, i32 undef>, <4 x i32> addrspace(3)* bitcast ([4 x <3 x i32>] addrspace(3)* @local_memory_alignment_global.lds_v3i32 to <4 x i32> addrspace(3)*), align 16
-// CHECK: store volatile <4 x i32> zeroinitializer, <4 x i32> addrspace(3)* getelementptr inbounds ([4 x <4 x i32>], [4 x <4 x i32>] addrspace(3)* @local_memory_alignment_global.lds_v4i32, i64 0, i64 0), align 16
-// CHECK: store volatile <8 x i32> zeroinitializer, <8 x i32> addrspace(3)* getelementptr inbounds ([4 x <8 x i32>], [4 x <8 x i32>] addrspace(3)* @local_memory_alignment_global.lds_v8i32, i64 0, i64 0), align 32
-// CHECK: store volatile <16 x i32> zeroinitializer, <16 x i32> addrspace(3)* getelementptr inbounds ([4 x <16 x i32>], [4 x <16 x i32>] addrspace(3)* @local_memory_alignment_global.lds_v16i32, i64 0, i64 0), align 64
-// CHECK: store volatile i64 0, i64 addrspace(3)* getelementptr inbounds ([4 x i64], [4 x i64] addrspace(3)* @local_memory_alignment_global.lds_i64, i64 0, i64 0), align 8
-// CHECK: store volatile <2 x i64> zeroinitializer, <2 x i64> addrspace(3)* getelementptr inbounds ([4 x <2 x i64>], [4 x <2 x i64>] addrspace(3)* @local_memory_alignment_global.lds_v2i64, i64 0, i64 0), align 16
-// CHECK: store volatile <4 x i64> <i64 0, i64 0, i64 0, i64 undef>, <4 x i64> addrspace(3)* bitcast ([4 x <3 x i64>] addrspace(3)* @local_memory_alignment_global.lds_v3i64 to <4 x i64> addrspace(3)*), align 32
-// CHECK: store volatile <4 x i64> zeroinitializer, <4 x i64> addrspace(3)* getelementptr inbounds ([4 x <4 x i64>], [4 x <4 x i64>] addrspace(3)* @local_memory_alignment_global.lds_v4i64, i64 0, i64 0), align 32
-// CHECK: store volatile <8 x i64> zeroinitializer, <8 x i64> addrspace(3)* getelementptr inbounds ([4 x <8 x i64>], [4 x <8 x i64>] addrspace(3)* @local_memory_alignment_global.lds_v8i64, i64 0, i64 0), align 64
-// CHECK: store volatile <16 x i64> zeroinitializer, <16 x i64> addrspace(3)* getelementptr inbounds ([4 x <16 x i64>], [4 x <16 x i64>] addrspace(3)* @local_memory_alignment_global.lds_v16i64, i64 0, i64 0), align 128
-// CHECK: store volatile half 0xH0000, half addrspace(3)* getelementptr inbounds ([4 x half], [4 x half] addrspace(3)* @local_memory_alignment_global.lds_f16, i64 0, i64 0), align 2
-// CHECK: store volatile <2 x half> zeroinitializer, <2 x half> addrspace(3)* getelementptr inbounds ([4 x <2 x half>], [4 x <2 x half>] addrspace(3)* @local_memory_alignment_global.lds_v2f16, i64 0, i64 0), align 4
-// CHECK: store volatile <4 x half> <half 0xH0000, half 0xH0000, half 0xH0000, half undef>, <4 x half> addrspace(3)* bitcast ([4 x <3 x half>] addrspace(3)* @local_memory_alignment_global.lds_v3f16 to <4 x half> addrspace(3)*), align 8
-// CHECK: store volatile <4 x half> zeroinitializer, <4 x half> addrspace(3)* getelementptr inbounds ([4 x <4 x half>], [4 x <4 x half>] addrspace(3)* @local_memory_alignment_global.lds_v4f16, i64 0, i64 0), align 8
-// CHECK: store volatile <8 x half> zeroinitializer, <8 x half> addrspace(3)* getelementptr inbounds ([4 x <8 x half>], [4 x <8 x half>] addrspace(3)* @local_memory_alignment_global.lds_v8f16, i64 0, i64 0), align 16
-// CHECK: store volatile <16 x half> zeroinitializer, <16 x half> addrspace(3)* getelementptr inbounds ([4 x <16 x half>], [4 x <16 x half>] addrspace(3)* @local_memory_alignment_global.lds_v16f16, i64 0, i64 0), align 32
-// CHECK: store volatile float 0.000000e+00, float addrspace(3)* getelementptr inbounds ([4 x float], [4 x float] addrspace(3)* @local_memory_alignment_global.lds_f32, i64 0, i64 0), align 4
-// CHECK: store volatile <2 x float> zeroinitializer, <2 x float> addrspace(3)* getelementptr inbounds ([4 x <2 x float>], [4 x <2 x float>] addrspace(3)* @local_memory_alignment_global.lds_v2f32, i64 0, i64 0), align 8
-// CHECK: store volatile <4 x float> <float 0.000000e+00, float 0.000000e+00, float 0.000000e+00, float undef>, <4 x float> addrspace(3)* bitcast ([4 x <3 x float>] addrspace(3)* @local_memory_alignment_global.lds_v3f32 to <4 x float> addrspace(3)*), align 16
-// CHECK: store volatile <4 x float> zeroinitializer, <4 x float> addrspace(3)* getelementptr inbounds ([4 x <4 x float>], [4 x <4 x float>] addrspace(3)* @local_memory_alignment_global.lds_v4f32, i64 0, i64 0), align 16
-// CHECK: store volatile <8 x float> zeroinitializer, <8 x float> addrspace(3)* getelementptr inbounds ([4 x <8 x float>], [4 x <8 x float>] addrspace(3)* @local_memory_alignment_global.lds_v8f32, i64 0, i64 0), align 32
-// CHECK: store volatile <16 x float> zeroinitializer, <16 x float> addrspace(3)* getelementptr inbounds ([4 x <16 x float>], [4 x <16 x float>] addrspace(3)* @local_memory_alignment_global.lds_v16f32, i64 0, i64 0), align 64
-// CHECK: store volatile double 0.000000e+00, double addrspace(3)* getelementptr inbounds ([4 x double], [4 x double] addrspace(3)* @local_memory_alignment_global.lds_f64, i64 0, i64 0), align 8
-// CHECK: store volatile <2 x double> zeroinitializer, <2 x double> addrspace(3)* getelementptr inbounds ([4 x <2 x double>], [4 x <2 x double>] addrspace(3)* @local_memory_alignment_global.lds_v2f64, i64 0, i64 0), align 16
-// CHECK: store volatile <4 x double> <double 0.000000e+00, double 0.000000e+00, double 0.000000e+00, double undef>, <4 x double> addrspace(3)* bitcast ([4 x <3 x double>] addrspace(3)* @local_memory_alignment_global.lds_v3f64 to <4 x double> addrspace(3)*), align 32
-// CHECK: store volatile <4 x double> zeroinitializer, <4 x double> addrspace(3)* getelementptr inbounds ([4 x <4 x double>], [4 x <4 x double>] addrspace(3)* @local_memory_alignment_global.lds_v4f64, i64 0, i64 0), align 32
-// CHECK: store volatile <8 x double> zeroinitializer, <8 x double> addrspace(3)* getelementptr inbounds ([4 x <8 x double>], [4 x <8 x double>] addrspace(3)* @local_memory_alignment_global.lds_v8f64, i64 0, i64 0), align 64
-// CHECK: store volatile <16 x double> zeroinitializer, <16 x double> addrspace(3)* getelementptr inbounds ([4 x <16 x double>], [4 x <16 x double>] addrspace(3)* @local_memory_alignment_global.lds_v16f64, i64 0, i64 0), align 128
+// CHECK: store volatile i8 0, ptr addrspace(3) @local_memory_alignment_global.lds_i8, align 1
+// CHECK: store volatile <2 x i8> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v2i8, align 2
+// CHECK: store volatile <4 x i8> <i8 0, i8 0, i8 0, i8 undef>, ptr addrspace(3) @local_memory_alignment_global.lds_v3i8, align 4
+// CHECK: store volatile <4 x i8> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v4i8, align 4
+// CHECK: store volatile <8 x i8> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v8i8, align 8
+// CHECK: store volatile <16 x i8> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v16i8, align 16
+// CHECK: store volatile i16 0, ptr addrspace(3) @local_memory_alignment_global.lds_i16, align 2
+// CHECK: store volatile <2 x i16> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v2i16, align 4
+// CHECK: store volatile <4 x i16> <i16 0, i16 0, i16 0, i16 undef>, ptr addrspace(3) @local_memory_alignment_global.lds_v3i16, align 8
+// CHECK: store volatile <4 x i16> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v4i16, align 8
+// CHECK: store volatile <8 x i16> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v8i16, align 16
+// CHECK: store volatile <16 x i16> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v16i16, align 32
+// CHECK: store volatile i32 0, ptr addrspace(3) @local_memory_alignment_global.lds_i32, align 4
+// CHECK: store volatile <2 x i32> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v2i32, align 8
+// CHECK: store volatile <4 x i32> <i32 0, i32 0, i32 0, i32 undef>, ptr addrspace(3) @local_memory_alignment_global.lds_v3i32, align 16
+// CHECK: store volatile <4 x i32> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v4i32, align 16
+// CHECK: store volatile <8 x i32> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v8i32, align 32
+// CHECK: store volatile <16 x i32> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v16i32, align 64
+// CHECK: store volatile i64 0, ptr addrspace(3) @local_memory_alignment_global.lds_i64, align 8
+// CHECK: store volatile <2 x i64> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v2i64, align 16
+// CHECK: store volatile <4 x i64> <i64 0, i64 0, i64 0, i64 undef>, ptr addrspace(3) @local_memory_alignment_global.lds_v3i64, align 32
+// CHECK: store volatile <4 x i64> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v4i64, align 32
+// CHECK: store volatile <8 x i64> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v8i64, align 64
+// CHECK: store volatile <16 x i64> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v16i64, align 128
+// CHECK: store volatile half 0xH0000, ptr addrspace(3) @local_memory_alignment_global.lds_f16, align 2
+// CHECK: store volatile <2 x half> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v2f16, align 4
+// CHECK: store volatile <4 x half> <half 0xH0000, half 0xH0000, half 0xH0000, half undef>, ptr addrspace(3) @local_memory_alignment_global.lds_v3f16, align 8
+// CHECK: store volatile <4 x half> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v4f16, align 8
+// CHECK: store volatile <8 x half> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v8f16, align 16
+// CHECK: store volatile <16 x half> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v16f16, align 32
+// CHECK: store volatile float 0.000000e+00, ptr addrspace(3) @local_memory_alignment_global.lds_f32, align 4
+// CHECK: store volatile <2 x float> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v2f32, align 8
+// CHECK: store volatile <4 x float> <float 0.000000e+00, float 0.000000e+00, float 0.000000e+00, float undef>, ptr addrspace(3) @local_memory_alignment_global.lds_v3f32, align 16
+// CHECK: store volatile <4 x float> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v4f32, align 16
+// CHECK: store volatile <8 x float> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v8f32, align 32
+// CHECK: store volatile <16 x float> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v16f32, align 64
+// CHECK: store volatile double 0.000000e+00, ptr addrspace(3) @local_memory_alignment_global.lds_f64, align 8
+// CHECK: store volatile <2 x double> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v2f64, align 16
+// CHECK: store volatile <4 x double> <double 0.000000e+00, double 0.000000e+00, double 0.000000e+00, double undef>, ptr addrspace(3) @local_memory_alignment_global.lds_v3f64, align 32
+// CHECK: store volatile <4 x double> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v4f64, align 32
+// CHECK: store volatile <8 x double> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v8f64, align 64
+// CHECK: store volatile <16 x double> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v16f64, align 128
 kernel void local_memory_alignment_global()
 {
   volatile local char lds_i8[4];
   volatile local char2 lds_v2i8[4];
   volatile local char3 lds_v3i8[4];
   volatile local char4 lds_v4i8[4];
   volatile local char8 lds_v8i8[4];
   volatile local char16 lds_v16i8[4];

 // CHECK: %private_v16f32 = alloca [4 x <16 x float>], align 64, addrspace(5)
 // CHECK: %private_f64 = alloca [4 x double], align 8, addrspace(5)
 // CHECK: %private_v2f64 = alloca [4 x <2 x double>], align 16, addrspace(5)
 // CHECK: %private_v3f64 = alloca [4 x <3 x double>], align 32, addrspace(5)
 // CHECK: %private_v4f64 = alloca [4 x <4 x double>], align 32, addrspace(5)
 // CHECK: %private_v8f64 = alloca [4 x <8 x double>], align 64, addrspace(5)
 // CHECK: %private_v16f64 = alloca [4 x <16 x double>], align 128, addrspace(5)
 
-// CHECK: store volatile i8 0, i8 addrspace(5)* %arraydecay, align 1
-// CHECK: store volatile <2 x i8> zeroinitializer, <2 x i8> addrspace(5)* %arraydecay{{[0-9]+}}, align 2
-// CHECK: store volatile <4 x i8> <i8 0, i8 0, i8 0, i8 undef>, <4 x i8> addrspace(5)* %storetmp, align 4
-// CHECK: store volatile <4 x i8> zeroinitializer, <4 x i8> addrspace(5)* %arraydecay{{[0-9]+}}, align 4
-// CHECK: store volatile <8 x i8> zeroinitializer, <8 x i8> addrspace(5)* %arraydecay{{[0-9]+}}, align 8
-// CHECK: store volatile <16 x i8> zeroinitializer, <16 x i8> addrspace(5)* %arraydecay{{[0-9]+}}, align 16
-// CHECK: store volatile i16 0, i16 addrspace(5)* %arraydecay{{[0-9]+}}, align 2
-// CHECK: store volatile <2 x i16> zeroinitializer, <2 x i16> addrspace(5)* %arraydecay{{[0-9]+}}, align 4
-// CHECK: store volatile <4 x i16> <i16 0, i16 0, i16 0, i16 undef>, <4 x i16> addrspace(5)* %storetmp{{[0-9]+}}, align 8
-// CHECK: store volatile <4 x i16> zeroinitializer, <4 x i16> addrspace(5)* %arraydecay{{[0-9]+}}, align 8
-// CHECK: store volatile <8 x i16> zeroinitializer, <8 x i16> addrspace(5)* %arraydecay{{[0-9]+}}, align 16
-// CHECK: store volatile <16 x i16> zeroinitializer, <16 x i16> addrspace(5)* %arraydecay{{[0-9]+}}, align 32
-// CHECK: store volatile i32 0, i32 addrspace(5)* %arraydecay{{[0-9]+}}, align 4
-// CHECK: store volatile <2 x i32> zeroinitializer, <2 x i32> addrspace(5)* %arraydecay{{[0-9]+}}, align 8
-// CHECK: store volatile <4 x i32> <i32 0, i32 0, i32 0, i32 undef>, <4 x i32> addrspace(5)* %storetmp16, align 16
-// CHECK: store volatile <4 x i32> zeroinitializer, <4 x i32> addrspace(5)* %arraydecay{{[0-9]+}}, align 16
-// CHECK: store volatile <8 x i32> zeroinitializer, <8 x i32> addrspace(5)* %arraydecay{{[0-9]+}}, align 32
-// CHECK: store volatile <16 x i32> zeroinitializer, <16 x i32> addrspace(5)* %arraydecay{{[0-9]+}}, align 64
-// CHECK: store volatile i64 0, i64 addrspace(5)* %arraydecay{{[0-9]+}}, align 8
-// CHECK: store volatile <2 x i64> zeroinitializer, <2 x i64> addrspace(5)* %arraydecay{{[0-9]+}}, align 16
-// CHECK: store volatile <4 x i64> <i64 0, i64 0, i64 0, i64 undef>, <4 x i64> addrspace(5)* %storetmp23, align 32
-// CHECK: store volatile <4 x i64> zeroinitializer, <4 x i64> addrspace(5)* %arraydecay{{[0-9]+}}, align 32
-// CHECK: store volatile <8 x i64> zeroinitializer, <8 x i64> addrspace(5)* %arraydecay{{[0-9]+}}, align 64
-// CHECK: store volatile <16 x i64> zeroinitializer, <16 x i64> addrspace(5)* %arraydecay{{[0-9]+}}, align 128
-// CHECK: store volatile half 0xH0000, half addrspace(5)* %arraydecay{{[0-9]+}}, align 2
-// CHECK: store volatile <2 x half> zeroinitializer, <2 x half> addrspace(5)* %arraydecay{{[0-9]+}}, align 4
-// CHECK: store volatile <4 x half> <half 0xH0000, half 0xH0000, half 0xH0000, half undef>, <4 x half> addrspace(5)* %storetmp{{[0-9]+}}, align 8
-// CHECK: store volatile <4 x half> zeroinitializer, <4 x half> addrspace(5)* %arraydecay{{[0-9]+}}, align 8
-// CHECK: store volatile <8 x half> zeroinitializer, <8 x half> addrspace(5)* %arraydecay{{[0-9]+}}, align 16
-// CHECK: store volatile <16 x half> zeroinitializer, <16 x half> addrspace(5)* %arraydecay{{[0-9]+}}, align 32
-// CHECK: store volatile float 0.000000e+00, float addrspace(5)* %arraydecay34, align 4
-// CHECK: store volatile <2 x float> zeroinitializer, <2 x float> addrspace(5)* %arraydecay{{[0-9]+}}, align 8
-// CHECK: store volatile <4 x float> <float 0.000000e+00, float 0.000000e+00, float 0.000000e+00, float undef>, <4 x float> addrspace(5)* %storetmp{{[0-9]+}}, align 16
-// CHECK: store volatile <4 x float> zeroinitializer, <4 x float> addrspace(5)* %arraydecay{{[0-9]+}}, align 16
-// CHECK: store volatile <8 x float> zeroinitializer, <8 x float> addrspace(5)* %arraydecay{{[0-9]+}}, align 32
-// CHECK: store volatile <16 x float> zeroinitializer, <16 x float> addrspace(5)* %arraydecay{{[0-9]+}}, align 64
-// CHECK: store volatile double 0.000000e+00, double addrspace(5)* %arraydecay{{[0-9]+}}, align 8
-// CHECK: store volatile <2 x double> zeroinitializer, <2 x double> addrspace(5)* %arraydecay{{[0-9]+}}, align 16
-// CHECK: store volatile <4 x double> <double 0.000000e+00, double 0.000000e+00, double 0.000000e+00, double undef>, <4 x double> addrspace(5)* %storetmp{{[0-9]+}}, align 32
-// CHECK: store volatile <4 x double> zeroinitializer, <4 x double> addrspace(5)* %arraydecay{{[0-9]+}}, align 32
-// CHECK: store volatile <8 x double> zeroinitializer, <8 x double> addrspace(5)* %arraydecay{{[0-9]+}}, align 64
-// CHECK: store volatile <16 x double> zeroinitializer, <16 x double> addrspace(5)* %arraydecay{{[0-9]+}}, align 128
+// CHECK: store volatile i8 0, ptr addrspace(5) %arraydecay, align 1
+// CHECK: store volatile <2 x i8> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 2
+// CHECK: store volatile <4 x i8> <i8 0, i8 0, i8 0, i8 undef>, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 4
+// CHECK: store volatile <4 x i8> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 4
+// CHECK: store volatile <8 x i8> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 8
+// CHECK: store volatile <16 x i8> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 16
+// CHECK: store volatile i16 0, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 2
+// CHECK: store volatile <2 x i16> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 4
+// CHECK: store volatile <4 x i16> <i16 0, i16 0, i16 0, i16 undef>, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 8
+// CHECK: store volatile <4 x i16> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 8
+// CHECK: store volatile <8 x i16> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 16
+// CHECK: store volatile <16 x i16> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 32
+// CHECK: store volatile i32 0, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 4
+// CHECK: store volatile <2 x i32> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 8
+// CHECK: store volatile <4 x i32> <i32 0, i32 0, i32 0, i32 undef>, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 16
+// CHECK: store volatile <4 x i32> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 16
+// CHECK: store volatile <8 x i32> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 32
+// CHECK: store volatile <16 x i32> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 64
+// CHECK: store volatile i64 0, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 8
+// CHECK: store volatile <2 x i64> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 16
+// CHECK: store volatile <4 x i64> <i64 0, i64 0, i64 0, i64 undef>, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 32
+// CHECK: store volatile <4 x i64> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 32
+// CHECK: store volatile <8 x i64> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 64
+// CHECK: store volatile <16 x i64> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 128
+// CHECK: store volatile half 0xH0000, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 2
+// CHECK: store volatile <2 x half> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 4
+// CHECK: store volatile <4 x half> <half 0xH0000, half 0xH0000, half 0xH0000, half undef>, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 8
+// CHECK: store volatile <4 x half> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 8
+// CHECK: store volatile <8 x half> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 16
+// CHECK: store volatile <16 x half> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 32
+// CHECK: store volatile float 0.000000e+00, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 4
+// CHECK: store volatile <2 x float> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 8
+// CHECK: store volatile <4 x float> <float 0.000000e+00, float 0.000000e+00, float 0.000000e+00, float undef>, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 16
+// CHECK: store volatile <4 x float> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 16
+// CHECK: store volatile <8 x float> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 32
+// CHECK: store volatile <16 x float> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 64
+// CHECK: store volatile double 0.000000e+00, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 8
+// CHECK: store volatile <2 x double> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 16
+// CHECK: store volatile <4 x double> <double 0.000000e+00, double 0.000000e+00, double 0.000000e+00, double undef>, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 32
+// CHECK: store volatile <4 x double> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 32
+// CHECK: store volatile <8 x double> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 64
+// CHECK: store volatile <16 x double> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 128
 kernel void private_memory_alignment_alloca()
 {
   volatile private char private_i8[4];
   volatile private char2 private_v2i8[4];
   volatile private char3 private_v3i8[4];
   volatile private char4 private_v4i8[4];
   volatile private char8 private_v8i8[4];
   volatile private char16 private_v16i8[4];

clang/test/CodeGenOpenCL/atomic-ops-libcall.cl

-// RUN: %clang_cc1 -no-opaque-pointers < %s -cl-std=CL2.0 -triple spir64 -emit-llvm | FileCheck -check-prefix=SPIR %s
-// RUN: %clang_cc1 -no-opaque-pointers < %s -cl-std=CL2.0 -triple armv5e-none-linux-gnueabi -emit-llvm | FileCheck -check-prefix=ARM %s
+// RUN: %clang_cc1 < %s -cl-std=CL2.0 -triple spir64 -emit-llvm | FileCheck -check-prefix=SPIR %s
+// RUN: %clang_cc1 < %s -cl-std=CL2.0 -triple armv5e-none-linux-gnueabi -emit-llvm | FileCheck -check-prefix=ARM %s
 typedef enum memory_order {
   memory_order_relaxed = __ATOMIC_RELAXED,
   memory_order_acquire = __ATOMIC_ACQUIRE,
   memory_order_release = __ATOMIC_RELEASE,
   memory_order_acq_rel = __ATOMIC_ACQ_REL,
   memory_order_seq_cst = __ATOMIC_SEQ_CST
 } memory_order;
 
 typedef enum memory_scope {
   memory_scope_work_item = __OPENCL_MEMORY_SCOPE_WORK_ITEM,
   memory_scope_work_group = __OPENCL_MEMORY_SCOPE_WORK_GROUP,
   memory_scope_device = __OPENCL_MEMORY_SCOPE_DEVICE,
   memory_scope_all_svm_devices = __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES,
 #if defined(cl_intel_subgroups) || defined(cl_khr_subgroups)
   memory_scope_sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP
 #endif
 } memory_scope;
 
 void f(atomic_int *i, global atomic_int *gi, local atomic_int *li, private atomic_int *pi, atomic_uint *ui, int cmp, int order, int scope) {
   int x;
-  // SPIR: {{%[^ ]*}} = call i32 @__opencl_atomic_load_4(i8 addrspace(4)* noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
-  // ARM: {{%[^ ]*}} = call i32 @__opencl_atomic_load_4(i8* noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
+  // SPIR: {{%[^ ]*}} = call i32 @__opencl_atomic_load_4(ptr addrspace(4) noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
+  // ARM: {{%[^ ]*}} = call i32 @__opencl_atomic_load_4(ptr noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
   x = __opencl_atomic_load(i, memory_order_seq_cst, memory_scope_work_group);
 
-  // SPIR: call void @__opencl_atomic_store_4(i8 addrspace(4)* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
-  // ARM: call void @__opencl_atomic_store_4(i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
+  // SPIR: call void @__opencl_atomic_store_4(ptr addrspace(4) noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
+  // ARM: call void @__opencl_atomic_store_4(ptr noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
   __opencl_atomic_store(i, 1, memory_order_seq_cst, memory_scope_work_group);
 
-  // SPIR: %[[GP:[0-9]+]] = addrspacecast i8 addrspace(1)* {{%[0-9]+}} to i8 addrspace(4)*
-  // SPIR: call void @__opencl_atomic_store_4(i8 addrspace(4)* noundef %[[GP]], i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
-  // ARM: call void @__opencl_atomic_store_4(i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
+  // SPIR: %[[GP:[0-9]+]] = addrspacecast ptr addrspace(1) {{%[0-9]+}} to ptr addrspace(4)
+  // SPIR: call void @__opencl_atomic_store_4(ptr addrspace(4) noundef %[[GP]], i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
+  // ARM: call void @__opencl_atomic_store_4(ptr noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
   __opencl_atomic_store(gi, 1, memory_order_seq_cst, memory_scope_work_group);
 
-  // SPIR: %[[GP:[0-9]+]] = addrspacecast i8 addrspace(3)* {{%[0-9]+}} to i8 addrspace(4)*
-  // SPIR: call void @__opencl_atomic_store_4(i8 addrspace(4)* noundef %[[GP]], i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
-  // ARM: call void @__opencl_atomic_store_4(i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
+  // SPIR: %[[GP:[0-9]+]] = addrspacecast ptr addrspace(3) {{%[0-9]+}} to ptr addrspace(4)
+  // SPIR: call void @__opencl_atomic_store_4(ptr addrspace(4) noundef %[[GP]], i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
+  // ARM: call void @__opencl_atomic_store_4(ptr noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
   __opencl_atomic_store(li, 1, memory_order_seq_cst, memory_scope_work_group);
 
-  // SPIR: %[[GP:[0-9]+]] = addrspacecast i8* {{%[0-9]+}} to i8 addrspace(4)*
-  // SPIR: call void @__opencl_atomic_store_4(i8 addrspace(4)* noundef %[[GP]], i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
-  // ARM: call void @__opencl_atomic_store_4(i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
+  // SPIR: %[[GP:[0-9]+]] = addrspacecast ptr {{%[0-9]+}} to ptr addrspace(4)
+  // SPIR: call void @__opencl_atomic_store_4(ptr addrspace(4) noundef %[[GP]], i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
+  // ARM: call void @__opencl_atomic_store_4(ptr noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
   __opencl_atomic_store(pi, 1, memory_order_seq_cst, memory_scope_work_group);
 
-  // SPIR: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_add_4(i8 addrspace(4)* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
-  // ARM: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_add_4(i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
+  // SPIR: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_add_4(ptr addrspace(4) noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
+  // ARM: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_add_4(ptr noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
   x = __opencl_atomic_fetch_add(i, 3, memory_order_seq_cst, memory_scope_work_group);
 
-  // SPIR: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_min_4(i8 addrspace(4)* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
-  // ARM: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_min_4(i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
+  // SPIR: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_min_4(ptr addrspace(4) noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
+  // ARM: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_min_4(ptr noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
   x = __opencl_atomic_fetch_min(i, 3, memory_order_seq_cst, memory_scope_work_group);
 
-  // SPIR: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_umin_4(i8 addrspace(4)* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
-  // ARM: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_umin_4(i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
+  // SPIR: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_umin_4(ptr addrspace(4) noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
+  // ARM: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_umin_4(ptr noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
   x = __opencl_atomic_fetch_min(ui, 3, memory_order_seq_cst, memory_scope_work_group);
 
-  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* noundef {{%[0-9]+}}, i8 addrspace(4)* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 1)
-  // ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8* noundef {{%[0-9]+}}, i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 1)
+  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(ptr addrspace(4) noundef {{%[0-9]+}}, ptr addrspace(4) noundef {{%[^,]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 1)
+  // ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(ptr noundef {{%[0-9]+}}, ptr noundef {{%[^,]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 1)
   x = __opencl_atomic_compare_exchange_strong(i, &cmp, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
 
-  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* noundef {{%[0-9]+}}, i8 addrspace(4)* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 1)
-  // ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8* noundef {{%[0-9]+}}, i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 1)
+  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(ptr addrspace(4) noundef {{%[0-9]+}}, ptr addrspace(4) noundef {{%[^,]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 1)
+  // ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(ptr noundef {{%[0-9]+}}, ptr noundef {{%[^,]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 1)
   x = __opencl_atomic_compare_exchange_weak(i, &cmp, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
 
-  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* noundef {{%[0-9]+}}, i8 addrspace(4)* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 2)
-  // ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8* noundef {{%[0-9]+}}, i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 2)
+  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(ptr addrspace(4) noundef {{%[0-9]+}}, ptr addrspace(4) noundef {{%[^,]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 2)
+  // ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(ptr noundef {{%[0-9]+}}, ptr noundef {{%[^,]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 2)
   x = __opencl_atomic_compare_exchange_weak(i, &cmp, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_device);
 
-  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* noundef {{%[0-9]+}}, i8 addrspace(4)* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 3)
-  // ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8* noundef {{%[0-9]+}}, i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 3)
+  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(ptr addrspace(4) noundef {{%[0-9]+}}, ptr addrspace(4) noundef {{%[^,]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 3)
+  // ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(ptr noundef {{%[0-9]+}}, ptr noundef {{%[^,]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 3)
   x = __opencl_atomic_compare_exchange_weak(i, &cmp, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_all_svm_devices);
 
 #ifdef cl_khr_subgroups
-  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* noundef {{%[0-9]+}}, i8 addrspace(4)* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 4)
+  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(ptr addrspace(4) noundef {{%[0-9]+}}, ptr addrspace(4) noundef {{%[^,]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 4)
   x = __opencl_atomic_compare_exchange_weak(i, &cmp, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_sub_group);
 #endif
 
-  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* noundef {{%[0-9]+}}, i8 addrspace(4)* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef %{{.*}}, i32 noundef %{{.*}}, i32 noundef %{{.*}})
-  // ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8* noundef {{%[0-9]+}}, i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef %{{.*}}, i32 noundef %{{.*}}, i32 noundef %{{.*}})
+  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(ptr addrspace(4) noundef {{%[0-9]+}}, ptr addrspace(4) noundef {{%[^,]+}}, i32 noundef {{%[0-9]+}}, i32 noundef %{{.*}}, i32 noundef %{{.*}}, i32 noundef %{{.*}})
+  // ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(ptr noundef {{%[0-9]+}}, ptr noundef {{%[^,]+}}, i32 noundef {{%[0-9]+}}, i32 noundef %{{.*}}, i32 noundef %{{.*}}, i32 noundef %{{.*}})
   x = __opencl_atomic_compare_exchange_weak(i, &cmp, 1, order, order, scope);
 }

clang/test/CodeGenOpenCL/atomic-ops.cl

-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL2.0 -emit-llvm -O0 -o - -triple=amdgcn-amd-amdhsa \
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -O0 -o - -triple=amdgcn-amd-amdhsa \
 // RUN:   | FileCheck %s
 
 // Also test serialization of atomic operations here, to avoid duplicating the test.
-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL2.0 -emit-pch -O0 -o %t -triple=amdgcn-amd-amdhsa
-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL2.0 -include-pch %t -O0 -triple=amdgcn-amd-amdhsa \
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-pch -O0 -o %t -triple=amdgcn-amd-amdhsa
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -include-pch %t -O0 -triple=amdgcn-amd-amdhsa \
 // RUN:   -emit-llvm -o - | FileCheck %s
 
 #ifndef ALREADY_INCLUDED
 #define ALREADY_INCLUDED
 
 #pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
 #pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable
 
#if defined(cl_intel_subgroups) || defined(cl_khr_subgroups)
   memory_scope_sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP
 #endif
 } memory_scope;
 
 atomic_int j;
 
 void fi1(atomic_int *i) {
   // CHECK-LABEL: @fi1
-  // CHECK: load atomic i32, i32* %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
+  // CHECK: load atomic i32, ptr %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
   int x = __opencl_atomic_load(i, memory_order_seq_cst, memory_scope_work_group);
 
-  // CHECK: load atomic i32, i32* %{{[.0-9A-Z_a-z]+}} syncscope("agent") seq_cst, align 4
+  // CHECK: load atomic i32, ptr %{{[.0-9A-Z_a-z]+}} syncscope("agent") seq_cst, align 4
   x = __opencl_atomic_load(i, memory_order_seq_cst, memory_scope_device);
 
-  // CHECK: load atomic i32, i32* %{{[.0-9A-Z_a-z]+}} seq_cst, align 4
+  // CHECK: load atomic i32, ptr %{{[.0-9A-Z_a-z]+}} seq_cst, align 4
   x = __opencl_atomic_load(i, memory_order_seq_cst, memory_scope_all_svm_devices);
 
-  // CHECK: load atomic i32, i32* %{{[.0-9A-Z_a-z]+}} syncscope("wavefront") seq_cst, align 4
+  // CHECK: load atomic i32, ptr %{{[.0-9A-Z_a-z]+}} syncscope("wavefront") seq_cst, align 4
   x = __opencl_atomic_load(i, memory_order_seq_cst, memory_scope_sub_group);
 }
 
 void fi2(atomic_int *i) {
   // CHECK-LABEL: @fi2
-  // CHECK: store atomic i32 %{{[.0-9A-Z_a-z]+}}, i32* %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
+  // CHECK: store atomic i32 %{{[.0-9A-Z_a-z]+}}, ptr %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
   __opencl_atomic_store(i, 1, memory_order_seq_cst, memory_scope_work_group);
 }
 
 void test_addr(global atomic_int *ig, private atomic_int *ip, local atomic_int *il) {
   // CHECK-LABEL: @test_addr
-  // CHECK: store atomic i32 %{{[.0-9A-Z_a-z]+}}, i32 addrspace(1)* %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
+  // CHECK: store atomic i32 %{{[.0-9A-Z_a-z]+}}, ptr addrspace(1) %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
   __opencl_atomic_store(ig, 1, memory_order_seq_cst, memory_scope_work_group);
 
-  // CHECK: store atomic i32 %{{[.0-9A-Z_a-z]+}}, i32 addrspace(5)* %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
+  // CHECK: store atomic i32 %{{[.0-9A-Z_a-z]+}}, ptr addrspace(5) %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
   __opencl_atomic_store(ip, 1, memory_order_seq_cst, memory_scope_work_group);
 
-  // CHECK: store atomic i32 %{{[.0-9A-Z_a-z]+}}, i32 addrspace(3)* %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
+  // CHECK: store atomic i32 %{{[.0-9A-Z_a-z]+}}, ptr addrspace(3) %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
   __opencl_atomic_store(il, 1, memory_order_seq_cst, memory_scope_work_group);
 }
 
 void fi3(atomic_int *i, atomic_uint *ui) {
   // CHECK-LABEL: @fi3
-  // CHECK: atomicrmw and i32* %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
+  // CHECK: atomicrmw and ptr %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
   int x = __opencl_atomic_fetch_and(i, 1, memory_order_seq_cst, memory_scope_work_group);
 
-  // CHECK: atomicrmw min i32* %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
+  // CHECK: atomicrmw min ptr %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
   x = __opencl_atomic_fetch_min(i, 1, memory_order_seq_cst, memory_scope_work_group);
 
-  // CHECK: atomicrmw max i32* %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
+  // CHECK: atomicrmw max ptr %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
   x = __opencl_atomic_fetch_max(i, 1, memory_order_seq_cst, memory_scope_work_group);
 
-  // CHECK: atomicrmw umin i32* %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
+  // CHECK: atomicrmw umin ptr %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
   x = __opencl_atomic_fetch_min(ui, 1, memory_order_seq_cst, memory_scope_work_group);
 
-  // CHECK: atomicrmw umax i32* %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
+  // CHECK: atomicrmw umax ptr %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
   x = __opencl_atomic_fetch_max(ui, 1, memory_order_seq_cst, memory_scope_work_group);
 }
 
 bool fi4(atomic_int *i) {
   // CHECK-LABEL: @fi4(
-  // CHECK: [[PAIR:%[.0-9A-Z_a-z]+]] = cmpxchg i32* [[PTR:%[.0-9A-Z_a-z]+]], i32 [[EXPECTED:%[.0-9A-Z_a-z]+]], i32 [[DESIRED:%[.0-9A-Z_a-z]+]] syncscope("workgroup-one-as") acquire acquire, align 4
+  // CHECK: [[PAIR:%[.0-9A-Z_a-z]+]] = cmpxchg ptr [[PTR:%[.0-9A-Z_a-z]+]], i32 [[EXPECTED:%[.0-9A-Z_a-z]+]], i32 [[DESIRED:%[.0-9A-Z_a-z]+]] syncscope("workgroup-one-as") acquire acquire, align 4
   // CHECK: [[OLD:%[.0-9A-Z_a-z]+]] = extractvalue { i32, i1 } [[PAIR]], 0
   // CHECK: [[CMP:%[.0-9A-Z_a-z]+]] = extractvalue { i32, i1 } [[PAIR]], 1
   // CHECK: br i1 [[CMP]], label %[[STORE_EXPECTED:[.0-9A-Z_a-z]+]], label %[[CONTINUE:[.0-9A-Z_a-z]+]]
   // CHECK: store i32 [[OLD]]
   int cmp = 0;
   return __opencl_atomic_compare_exchange_strong(i, &cmp, 1, memory_order_acquire, memory_order_acquire, memory_scope_work_group);
 }
 
 void fi5(atomic_int *i, int scope) {
   // CHECK-LABEL: @fi5
   // CHECK: switch i32 %{{.*}}, label %[[opencl_allsvmdevices:.*]] [
   // CHECK-NEXT: i32 1, label %[[opencl_workgroup:.*]]
   // CHECK-NEXT: i32 2, label %[[opencl_device:.*]]
   // CHECK-NEXT: i32 4, label %[[opencl_subgroup:.*]]
   // CHECK-NEXT: ]
   // CHECK: [[opencl_workgroup]]:
-  // CHECK: load atomic i32, i32* %{{.*}} syncscope("workgroup") seq_cst, align 4
+  // CHECK: load atomic i32, ptr %{{.*}} syncscope("workgroup") seq_cst, align 4
   // CHECK: br label %[[continue:.*]]
   // CHECK: [[opencl_device]]:
-  // CHECK: load atomic i32, i32* %{{.*}} syncscope("agent") seq_cst, align 4
+  // CHECK: load atomic i32, ptr %{{.*}} syncscope("agent") seq_cst, align 4
   // CHECK: br label %[[continue]]
   // CHECK: [[opencl_allsvmdevices]]:
-  // CHECK: load atomic i32, i32* %{{.*}} seq_cst, align 4
+  // CHECK: load atomic i32, ptr %{{.*}} seq_cst, align 4
   // CHECK: br label %[[continue]]
   // CHECK: [[opencl_subgroup]]:
-  // CHECK: load atomic i32, i32* %{{.*}} syncscope("wavefront") seq_cst, align 4
+  // CHECK: load atomic i32, ptr %{{.*}} syncscope("wavefront") seq_cst, align 4
   // CHECK: br label %[[continue]]
   // CHECK: [[continue]]:
   int x = __opencl_atomic_load(i, memory_order_seq_cst, scope);
 }
 
 void fi6(atomic_int *i, int order, int scope) {
   // CHECK-LABEL: @fi6
   // CHECK: switch i32 %{{.*}}, label %[[monotonic:.*]] [
void fi6(atomic_int *i, int order, int scope) {
   // CHECK-NEXT: ]
   // CHECK: [[seqcst]]:
   // CHECK: switch i32 %{{.*}}, label %[[SEQ_ALL:.*]] [
   // CHECK-NEXT: i32 1, label %[[SEQ_WG:.*]]
   // CHECK-NEXT: i32 2, label %[[SEQ_DEV:.*]]
   // CHECK-NEXT: i32 4, label %[[SEQ_SUB:.*]]
   // CHECK-NEXT: ]
   // CHECK: [[MON_WG]]:
-  // CHECK: load atomic i32, i32* %{{.*}} syncscope("workgroup-one-as") monotonic, align 4
+  // CHECK: load atomic i32, ptr %{{.*}} syncscope("workgroup-one-as") monotonic, align 4
   // CHECK: [[MON_DEV]]:
-  // CHECK: load atomic i32, i32* %{{.*}} syncscope("agent-one-as") monotonic, align 4
+  // CHECK: load atomic i32, ptr %{{.*}} syncscope("agent-one-as") monotonic, align 4
   // CHECK: [[MON_ALL]]:
-  // CHECK: load atomic i32, i32* %{{.*}} monotonic, align 4
+  // CHECK: load atomic i32, ptr %{{.*}} monotonic, align 4
   // CHECK: [[MON_SUB]]:
-  // CHECK: load atomic i32, i32* %{{.*}} syncscope("wavefront-one-as") monotonic, align 4
+  // CHECK: load atomic i32, ptr %{{.*}} syncscope("wavefront-one-as") monotonic, align 4
   // CHECK: [[ACQ_WG]]:
-  // CHECK: load atomic i32, i32* %{{.*}} syncscope("workgroup-one-as") acquire, align 4
+  // CHECK: load atomic i32, ptr %{{.*}} syncscope("workgroup-one-as") acquire, align 4
   // CHECK: [[ACQ_DEV]]:
-  // CHECK: load atomic i32, i32* %{{.*}} syncscope("agent-one-as") acquire, align 4
+  // CHECK: load atomic i32, ptr %{{.*}} syncscope("agent-one-as") acquire, align 4
   // CHECK: [[ACQ_ALL]]:
-  // CHECK: load atomic i32, i32* %{{.*}} acquire, align 4
+  // CHECK: load atomic i32, ptr %{{.*}} acquire, align 4
   // CHECK: [[ACQ_SUB]]:
-  // CHECK: load atomic i32, i32* %{{.*}} syncscope("wavefront-one-as") acquire, align 4
+  // CHECK: load atomic i32, ptr %{{.*}} syncscope("wavefront-one-as") acquire, align 4
   // CHECK: [[SEQ_WG]]:
-  // CHECK: load atomic i32, i32* %{{.*}} syncscope("workgroup") seq_cst, align 4
+  // CHECK: load atomic i32, ptr %{{.*}} syncscope("workgroup") seq_cst, align 4
   // CHECK: [[SEQ_DEV]]:
-  // CHECK: load atomic i32, i32* %{{.*}} syncscope("agent") seq_cst, align 4
+  // CHECK: load atomic i32, ptr %{{.*}} syncscope("agent") seq_cst, align 4
   // CHECK: [[SEQ_ALL]]:
-  // CHECK: load atomic i32, i32* %{{.*}} seq_cst, align 4
+  // CHECK: load atomic i32, ptr %{{.*}} seq_cst, align 4
   // CHECK: [[SEQ_SUB]]:
-  // CHECK: load atomic i32, i32* %{{.*}} syncscope("wavefront") seq_cst, align 4
+  // CHECK: load atomic i32, ptr %{{.*}} syncscope("wavefront") seq_cst, align 4
   int x = __opencl_atomic_load(i, order, scope);
 }
 
 float ff1(global atomic_float *d) {
   // CHECK-LABEL: @ff1
-  // CHECK: load atomic i32, i32 addrspace(1)* {{.*}} syncscope("workgroup-one-as") monotonic, align 4
+  // CHECK: load atomic i32, ptr addrspace(1) {{.*}} syncscope("workgroup-one-as") monotonic, align 4
   return __opencl_atomic_load(d, memory_order_relaxed, memory_scope_work_group);
 }
 
 void ff2(atomic_float *d) {
   // CHECK-LABEL: @ff2
   // CHECK: store atomic i32 {{.*}} syncscope("workgroup-one-as") release, align 4
   __opencl_atomic_store(d, 1, memory_order_release, memory_scope_work_group);
 }
 
 float ff3(atomic_float *d) {
   // CHECK-LABEL: @ff3
-  // CHECK: atomicrmw xchg i32* {{.*}} syncscope("workgroup") seq_cst, align 4
+  // CHECK: atomicrmw xchg ptr {{.*}} syncscope("workgroup") seq_cst, align 4
   return __opencl_atomic_exchange(d, 2, memory_order_seq_cst, memory_scope_work_group);
 }
 
 float ff4(global atomic_float *d, float a) {
   // CHECK-LABEL: @ff4
-  // CHECK: atomicrmw fadd float addrspace(1)* {{.*}} syncscope("workgroup-one-as") monotonic
+  // CHECK: atomicrmw fadd ptr addrspace(1) {{.*}} syncscope("workgroup-one-as") monotonic
   return __opencl_atomic_fetch_add(d, a, memory_order_relaxed, memory_scope_work_group);
 }
 
 float ff5(global atomic_double *d, double a) {
   // CHECK-LABEL: @ff5
-  // CHECK: atomicrmw fadd double addrspace(1)* {{.*}} syncscope("workgroup-one-as") monotonic
+  // CHECK: atomicrmw fadd ptr addrspace(1) {{.*}} syncscope("workgroup-one-as") monotonic
   return __opencl_atomic_fetch_add(d, a, memory_order_relaxed, memory_scope_work_group);
 }
 
 // CHECK-LABEL: @atomic_init_foo
 void atomic_init_foo()
 {
   // CHECK-NOT: atomic
   // CHECK: store
   __opencl_atomic_init(&j, 42);
 
   // CHECK-NOT: atomic
   // CHECK: }
 }
 
 // CHECK-LABEL: @failureOrder
 void failureOrder(atomic_int *ptr, int *ptr2) {
-  // CHECK: cmpxchg i32* {{%[0-9A-Za-z._]+}}, i32 {{%[0-9A-Za-z._]+}}, i32 {{%[0-9A-Za-z_.]+}} syncscope("workgroup-one-as") acquire monotonic, align 4
+  // CHECK: cmpxchg ptr {{%[0-9A-Za-z._]+}}, i32 {{%[0-9A-Za-z._]+}}, i32 {{%[0-9A-Za-z_.]+}} syncscope("workgroup-one-as") acquire monotonic, align 4
   __opencl_atomic_compare_exchange_strong(ptr, ptr2, 43, memory_order_acquire, memory_order_relaxed, memory_scope_work_group);
 
-  // CHECK: cmpxchg weak i32* {{%[0-9A-Za-z._]+}}, i32 {{%[0-9A-Za-z._]+}}, i32 {{%[0-9A-Za-z_.]+}} syncscope("workgroup") seq_cst acquire, align 4
+  // CHECK: cmpxchg weak ptr {{%[0-9A-Za-z._]+}}, i32 {{%[0-9A-Za-z._]+}}, i32 {{%[0-9A-Za-z_.]+}} syncscope("workgroup") seq_cst acquire, align 4
   __opencl_atomic_compare_exchange_weak(ptr, ptr2, 43, memory_order_seq_cst, memory_order_acquire, memory_scope_work_group);
 }
 
 // CHECK-LABEL: @generalFailureOrder
 void generalFailureOrder(atomic_int *ptr, int *ptr2, int success, int fail) {
   __opencl_atomic_compare_exchange_strong(ptr, ptr2, 42, success, fail, memory_scope_work_group);
 // CHECK: switch i32 {{.*}}, label %[[MONOTONIC:[0-9a-zA-Z._]+]] [
   // CHECK-NEXT: i32 1, label %[[ACQUIRE:[0-9a-zA-Z._]+]]
// CHECK: switch i32 {{.*}}, label %[[MONOTONIC:[0-9a-zA-Z._]+]] [
 
   // CHECK: [[SEQCST_SEQCST]]
   // CHECK: cmpxchg {{.*}} seq_cst seq_cst, align 4
   // CHECK: br
 }
 
 int test_volatile(volatile atomic_int *i) {
   // CHECK-LABEL: @test_volatile
-  // CHECK:      %[[i_addr:.*]] = alloca i32
+  // CHECK:      %[[i_addr:.*]] = alloca ptr
   // CHECK-NEXT: %[[atomicdst:.*]] = alloca i32
-  // CHECK-NEXT: store i32* %i, i32* addrspace(5)* %[[i_addr]]
-  // CHECK-NEXT: %[[addr:.*]] = load i32*, i32* addrspace(5)* %[[i_addr]]
-  // CHECK-NEXT: %[[res:.*]] = load atomic volatile i32, i32* %[[addr]] syncscope("workgroup") seq_cst, align 4
-  // CHECK-NEXT: store i32 %[[res]], i32 addrspace(5)* %[[atomicdst]]
-  // CHECK-NEXT: %[[retval:.*]] = load i32, i32 addrspace(5)* %[[atomicdst]]
+  // CHECK-NEXT: store ptr %i, ptr addrspace(5) %[[i_addr]]
+  // CHECK-NEXT: %[[addr:.*]] = load ptr, ptr addrspace(5) %[[i_addr]]
+  // CHECK-NEXT: %[[res:.*]] = load atomic volatile i32, ptr %[[addr]] syncscope("workgroup") seq_cst, align 4
+  // CHECK-NEXT: store i32 %[[res]], ptr addrspace(5) %[[atomicdst]]
+  // CHECK-NEXT: %[[retval:.*]] = load i32, ptr addrspace(5) %[[atomicdst]]
   // CHECK-NEXT: ret i32 %[[retval]]
   return __opencl_atomic_load(i, memory_order_seq_cst, memory_scope_work_group);
 }
 
 #endif

clang/test/CodeGenOpenCL/blocks.cl

-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL2.0 -emit-llvm -o - -O0 -triple spir-unknown-unknown | FileCheck -check-prefixes=COMMON,SPIR %s
-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL2.0 -emit-llvm -o - -O0 -triple amdgcn-amd-amdhsa | FileCheck -check-prefixes=COMMON,AMDGCN %s
-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL2.0 -emit-llvm -o - -O0 -debug-info-kind=limited -triple spir-unknown-unknown | FileCheck -check-prefixes=CHECK-DEBUG %s
-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL2.0 -emit-llvm -o - -O0 -debug-info-kind=limited -triple amdgcn-amd-amdhsa | FileCheck -check-prefixes=CHECK-DEBUG %s
-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL3.0 -cl-ext=-all,+__opencl_c_device_enqueue,+__opencl_c_generic_address_space,+__opencl_c_program_scope_global_variables -emit-llvm -o - -O0 -triple spir-unknown-unknown | FileCheck -check-prefixes=COMMON,SPIR %s
-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL3.0 -cl-ext=-all,+__opencl_c_device_enqueue,+__opencl_c_generic_address_space,+__opencl_c_program_scope_global_variables  -emit-llvm -o - -O0 -triple amdgcn-amd-amdhsa | FileCheck -check-prefixes=COMMON,AMDGCN %s
-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL3.0 -cl-ext=-all,+__opencl_c_device_enqueue,+__opencl_c_generic_address_space,+__opencl_c_program_scope_global_variables  -emit-llvm -o - -O0 -debug-info-kind=limited -triple spir-unknown-unknown | FileCheck -check-prefixes=CHECK-DEBUG %s
-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL3.0 -cl-ext=-all,+__opencl_c_device_enqueue,+__opencl_c_generic_address_space,+__opencl_c_program_scope_global_variables -emit-llvm -o - -O0 -debug-info-kind=limited -triple amdgcn-amd-amdhsa | FileCheck -check-prefixes=CHECK-DEBUG %s
-
-// SPIR: %struct.__opencl_block_literal_generic = type { i32, i32, i8 addrspace(4)* }
-// AMDGCN: %struct.__opencl_block_literal_generic = type { i32, i32, i8* }
-// SPIR: @__block_literal_global = internal addrspace(1) constant { i32, i32, i8 addrspace(4)* } { i32 12, i32 4, i8 addrspace(4)* addrspacecast (i8* bitcast (void (i8 addrspace(4)*, i8 addrspace(3)*)* @block_A_block_invoke to i8*) to i8 addrspace(4)*) }
-// AMDGCN: @__block_literal_global = internal addrspace(1) constant { i32, i32, i8* } { i32 16, i32 8, i8* bitcast (void (i8*, i8 addrspace(3)*)* @block_A_block_invoke to i8*) }
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -O0 -triple spir-unknown-unknown | FileCheck -check-prefixes=COMMON,SPIR %s
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -O0 -triple amdgcn-amd-amdhsa | FileCheck -check-prefixes=COMMON,AMDGCN %s
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -O0 -debug-info-kind=limited -triple spir-unknown-unknown | FileCheck -check-prefixes=CHECK-DEBUG %s
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -O0 -debug-info-kind=limited -triple amdgcn-amd-amdhsa | FileCheck -check-prefixes=CHECK-DEBUG %s
+// RUN: %clang_cc1 %s -cl-std=CL3.0 -cl-ext=-all,+__opencl_c_device_enqueue,+__opencl_c_generic_address_space,+__opencl_c_program_scope_global_variables -emit-llvm -o - -O0 -triple spir-unknown-unknown | FileCheck -check-prefixes=COMMON,SPIR %s
+// RUN: %clang_cc1 %s -cl-std=CL3.0 -cl-ext=-all,+__opencl_c_device_enqueue,+__opencl_c_generic_address_space,+__opencl_c_program_scope_global_variables  -emit-llvm -o - -O0 -triple amdgcn-amd-amdhsa | FileCheck -check-prefixes=COMMON,AMDGCN %s
+// RUN: %clang_cc1 %s -cl-std=CL3.0 -cl-ext=-all,+__opencl_c_device_enqueue,+__opencl_c_generic_address_space,+__opencl_c_program_scope_global_variables  -emit-llvm -o - -O0 -debug-info-kind=limited -triple spir-unknown-unknown | FileCheck -check-prefixes=CHECK-DEBUG %s
+// RUN: %clang_cc1 %s -cl-std=CL3.0 -cl-ext=-all,+__opencl_c_device_enqueue,+__opencl_c_generic_address_space,+__opencl_c_program_scope_global_variables -emit-llvm -o - -O0 -debug-info-kind=limited -triple amdgcn-amd-amdhsa | FileCheck -check-prefixes=CHECK-DEBUG %s
+
+// SPIR: @__block_literal_global = internal addrspace(1) constant { i32, i32, ptr addrspace(4) } { i32 12, i32 4, ptr addrspace(4) addrspacecast (ptr @block_A_block_invoke to ptr addrspace(4)) }
+// AMDGCN: @__block_literal_global = internal addrspace(1) constant { i32, i32, ptr } { i32 16, i32 8, ptr @block_A_block_invoke }
 // COMMON-NOT: .str
 
-// SPIR-LABEL: define internal {{.*}}void @block_A_block_invoke(i8 addrspace(4)* noundef %.block_descriptor, i8 addrspace(3)* noundef %a)
-// AMDGCN-LABEL: define internal {{.*}}void @block_A_block_invoke(i8* noundef %.block_descriptor, i8 addrspace(3)* noundef %a)
+// SPIR-LABEL: define internal {{.*}}void @block_A_block_invoke(ptr addrspace(4) noundef %.block_descriptor, ptr addrspace(3) noundef %a)
+// AMDGCN-LABEL: define internal {{.*}}void @block_A_block_invoke(ptr noundef %.block_descriptor, ptr addrspace(3) noundef %a)
 void (^block_A)(local void *) = ^(local void *a) {
   return;
 };
 
 // COMMON-LABEL: define {{.*}}void @foo()
 void foo(){
   int i;
   // COMMON-NOT: %block.isa
   // COMMON-NOT: %block.flags
   // COMMON-NOT: %block.reserved
   // COMMON-NOT: %block.descriptor
-  // SPIR: %[[block_size:.*]] = getelementptr inbounds <{ i32, i32, i8 addrspace(4)*, i32 }>, <{ i32, i32, i8 addrspace(4)*, i32 }>* %block, i32 0, i32 0
-  // AMDGCN: %[[block_size:.*]] = getelementptr inbounds <{ i32, i32, i8*, i32 }>, <{ i32, i32, i8*, i32 }> addrspace(5)* %block, i32 0, i32 0
-  // SPIR: store i32 16, i32* %[[block_size]]
-  // AMDGCN: store i32 20, i32 addrspace(5)* %[[block_size]]
-  // SPIR: %[[block_align:.*]] = getelementptr inbounds <{ i32, i32, i8 addrspace(4)*, i32 }>, <{ i32, i32, i8 addrspace(4)*, i32 }>* %block, i32 0, i32 1
-  // AMDGCN: %[[block_align:.*]] = getelementptr inbounds <{ i32, i32, i8*, i32 }>, <{ i32, i32, i8*, i32 }> addrspace(5)* %block, i32 0, i32 1
-  // SPIR: store i32 4, i32* %[[block_align]]
-  // AMDGCN: store i32 8, i32 addrspace(5)* %[[block_align]]
-  // SPIR: %[[block_invoke:.*]] = getelementptr inbounds <{ i32, i32, i8 addrspace(4)*, i32 }>, <{ i32, i32, i8 addrspace(4)*, i32 }>* %[[block:.*]], i32 0, i32 2
-  // SPIR: store i8 addrspace(4)* addrspacecast (i8* bitcast (i32 (i8 addrspace(4)*)* @__foo_block_invoke to i8*) to i8 addrspace(4)*), i8 addrspace(4)** %[[block_invoke]]
-  // SPIR: %[[block_captured:.*]] = getelementptr inbounds <{ i32, i32, i8 addrspace(4)*, i32 }>, <{ i32, i32, i8 addrspace(4)*, i32 }>* %[[block]], i32 0, i32 3
-  // SPIR: %[[i_value:.*]] = load i32, i32* %i
-  // SPIR: store i32 %[[i_value]], i32* %[[block_captured]],
-  // SPIR: %[[blk_ptr:.*]] = bitcast <{ i32, i32, i8 addrspace(4)*, i32 }>* %[[block]] to %struct.__opencl_block_literal_generic*
-  // SPIR: %[[blk_gen_ptr:.*]] = addrspacecast %struct.__opencl_block_literal_generic* %[[blk_ptr]] to %struct.__opencl_block_literal_generic addrspace(4)*
-  // SPIR: store %struct.__opencl_block_literal_generic addrspace(4)* %[[blk_gen_ptr]], %struct.__opencl_block_literal_generic addrspace(4)** %[[block_B:.*]],
-  // SPIR: %[[block_literal:.*]] = load %struct.__opencl_block_literal_generic addrspace(4)*, %struct.__opencl_block_literal_generic addrspace(4)** %[[block_B]]
-  // SPIR: %[[blk_gen_ptr:.*]] = bitcast %struct.__opencl_block_literal_generic addrspace(4)* %[[block_literal]] to i8 addrspace(4)*
-  // SPIR: call {{.*}}i32 @__foo_block_invoke(i8 addrspace(4)* noundef %[[blk_gen_ptr]])
-  // AMDGCN: %[[block_invoke:.*]] = getelementptr inbounds <{ i32, i32, i8*, i32 }>, <{ i32, i32, i8*, i32 }> addrspace(5)* %[[block:.*]], i32 0, i32 2
-  // AMDGCN: store i8* bitcast (i32 (i8*)* @__foo_block_invoke to i8*), i8* addrspace(5)* %[[block_invoke]]
-  // AMDGCN: %[[block_captured:.*]] = getelementptr inbounds <{ i32, i32, i8*, i32 }>, <{ i32, i32, i8*, i32 }> addrspace(5)* %[[block]], i32 0, i32 3
-  // AMDGCN: %[[i_value:.*]] = load i32, i32 addrspace(5)* %i
-  // AMDGCN: store i32 %[[i_value]], i32 addrspace(5)* %[[block_captured]],
-  // AMDGCN: %[[blk_ptr:.*]] = bitcast <{ i32, i32, i8*, i32 }> addrspace(5)* %[[block]] to %struct.__opencl_block_literal_generic addrspace(5)*
-  // AMDGCN: %[[blk_gen_ptr:.*]] = addrspacecast %struct.__opencl_block_literal_generic addrspace(5)* %[[blk_ptr]] to %struct.__opencl_block_literal_generic*
-  // AMDGCN: store %struct.__opencl_block_literal_generic* %[[blk_gen_ptr]], %struct.__opencl_block_literal_generic* addrspace(5)* %[[block_B:.*]],
-  // AMDGCN: %[[block_literal:.*]] = load %struct.__opencl_block_literal_generic*, %struct.__opencl_block_literal_generic* addrspace(5)* %[[block_B]]
-  // AMDGCN: %[[blk_gen_ptr:.*]] = bitcast %struct.__opencl_block_literal_generic* %[[block_literal]] to i8*
-  // AMDGCN: call {{.*}}i32 @__foo_block_invoke(i8* noundef %[[blk_gen_ptr]])
+  // SPIR: %[[block_size:.*]] = getelementptr inbounds <{ i32, i32, ptr addrspace(4), i32 }>, ptr %block, i32 0, i32 0
+  // AMDGCN: %[[block_size:.*]] = getelementptr inbounds <{ i32, i32, ptr, i32 }>, ptr addrspace(5) %block, i32 0, i32 0
+  // SPIR: store i32 16, ptr %[[block_size]]
+  // AMDGCN: store i32 20, ptr addrspace(5) %[[block_size]]
+  // SPIR: %[[block_align:.*]] = getelementptr inbounds <{ i32, i32, ptr addrspace(4), i32 }>, ptr %block, i32 0, i32 1
+  // AMDGCN: %[[block_align:.*]] = getelementptr inbounds <{ i32, i32, ptr, i32 }>, ptr addrspace(5) %block, i32 0, i32 1
+  // SPIR: store i32 4, ptr %[[block_align]]
+  // AMDGCN: store i32 8, ptr addrspace(5) %[[block_align]]
+  // SPIR: %[[block_invoke:.*]] = getelementptr inbounds <{ i32, i32, ptr addrspace(4), i32 }>, ptr %[[block:.*]], i32 0, i32 2
+  // SPIR: store ptr addrspace(4) addrspacecast (ptr @__foo_block_invoke to ptr addrspace(4)), ptr %[[block_invoke]]
+  // SPIR: %[[block_captured:.*]] = getelementptr inbounds <{ i32, i32, ptr addrspace(4), i32 }>, ptr %[[block]], i32 0, i32 3
+  // SPIR: %[[i_value:.*]] = load i32, ptr %i
+  // SPIR: store i32 %[[i_value]], ptr %[[block_captured]],
+  // SPIR: %[[blk_gen_ptr:.*]] = addrspacecast ptr %[[block]] to ptr addrspace(4)
+  // SPIR: store ptr addrspace(4) %[[blk_gen_ptr]], ptr %[[block_B:.*]],
+  // SPIR: %[[block_literal:.*]] = load ptr addrspace(4), ptr %[[block_B]]
+  // SPIR: call {{.*}}i32 @__foo_block_invoke(ptr addrspace(4) noundef %[[block_literal]])
+  // AMDGCN: %[[block_invoke:.*]] = getelementptr inbounds <{ i32, i32, ptr, i32 }>, ptr addrspace(5) %[[block:.*]], i32 0, i32 2
+  // AMDGCN: store ptr @__foo_block_invoke, ptr addrspace(5) %[[block_invoke]]
+  // AMDGCN: %[[block_captured:.*]] = getelementptr inbounds <{ i32, i32, ptr, i32 }>, ptr addrspace(5) %[[block]], i32 0, i32 3
+  // AMDGCN: %[[i_value:.*]] = load i32, ptr addrspace(5) %i
+  // AMDGCN: store i32 %[[i_value]], ptr addrspace(5) %[[block_captured]],
+  // AMDGCN: %[[blk_gen_ptr:.*]] = addrspacecast ptr addrspace(5) %[[block]] to ptr
+  // AMDGCN: store ptr %[[blk_gen_ptr]], ptr addrspace(5) %[[block_B:.*]],
+  // AMDGCN: %[[block_literal:.*]] = load ptr, ptr addrspace(5) %[[block_B]]
+  // AMDGCN: call {{.*}}i32 @__foo_block_invoke(ptr noundef %[[block_literal]])
 
   int (^ block_B)(void) = ^{
     return i;
   };
   block_B();
 }
 
-// SPIR-LABEL: define internal {{.*}}i32 @__foo_block_invoke(i8 addrspace(4)* noundef %.block_descriptor)
-// SPIR:  %[[block:.*]] = bitcast i8 addrspace(4)* %.block_descriptor to <{ i32, i32, i8 addrspace(4)*, i32 }> addrspace(4)*
-// SPIR:  %[[block_capture_addr:.*]] = getelementptr inbounds <{ i32, i32, i8 addrspace(4)*, i32 }>, <{ i32, i32, i8 addrspace(4)*, i32 }> addrspace(4)* %[[block]], i32 0, i32 3
-// SPIR:  %[[block_capture:.*]] = load i32, i32 addrspace(4)* %[[block_capture_addr]]
-// AMDGCN-LABEL: define internal {{.*}}i32 @__foo_block_invoke(i8* noundef %.block_descriptor)
-// AMDGCN:  %[[block:.*]] = bitcast i8* %.block_descriptor to <{ i32, i32, i8*, i32 }>*
-// AMDGCN:  %[[block_capture_addr:.*]] = getelementptr inbounds <{ i32, i32, i8*, i32 }>, <{ i32, i32, i8*, i32 }>* %[[block]], i32 0, i32 3
-// AMDGCN:  %[[block_capture:.*]] = load i32, i32* %[[block_capture_addr]]
+// SPIR-LABEL: define internal {{.*}}i32 @__foo_block_invoke(ptr addrspace(4) noundef %.block_descriptor)
+// SPIR:  %[[block_capture_addr:.*]] = getelementptr inbounds <{ i32, i32, ptr addrspace(4), i32 }>, ptr addrspace(4) %.block_descriptor, i32 0, i32 3
+// SPIR:  %[[block_capture:.*]] = load i32, ptr addrspace(4) %[[block_capture_addr]]
+// AMDGCN-LABEL: define internal {{.*}}i32 @__foo_block_invoke(ptr noundef %.block_descriptor)
+// AMDGCN:  %[[block_capture_addr:.*]] = getelementptr inbounds <{ i32, i32, ptr, i32 }>, ptr %.block_descriptor, i32 0, i32 3
+// AMDGCN:  %[[block_capture:.*]] = load i32, ptr %[[block_capture_addr]]
 
 // COMMON-NOT: define{{.*}}@__foo_block_invoke_kernel
 
 // COMMON-LABEL: define {{.*}}@call_block
 // call {{.*}}@__call_block_block_invoke
 int call_block() {
   return ^int(int num) { return num; } (11);
 }
 
 // CHECK-DEBUG: !DIDerivedType(tag: DW_TAG_member, name: "__size"
 // CHECK-DEBUG: !DIDerivedType(tag: DW_TAG_member, name: "__align"
 
 // CHECK-DEBUG-NOT: !DIDerivedType(tag: DW_TAG_member, name: "__isa"
 // CHECK-DEBUG-NOT: !DIDerivedType(tag: DW_TAG_member, name: "__flags"
 // CHECK-DEBUG-NOT: !DIDerivedType(tag: DW_TAG_member, name: "__reserved"
 // CHECK-DEBUG-NOT: !DIDerivedType(tag: DW_TAG_member, name: "__FuncPtr"

clang/test/CodeGenOpenCL/builtins.cl

-// RUN: %clang_cc1 -no-opaque-pointers %s -finclude-default-header -fdeclare-opencl-builtins -cl-std=clc++ -fblocks -O0 -emit-llvm -o - -triple "spir-unknown-unknown" | FileCheck %s
+// RUN: %clang_cc1 %s -finclude-default-header -fdeclare-opencl-builtins -cl-std=clc++ -fblocks -O0 -emit-llvm -o - -triple "spir-unknown-unknown" | FileCheck %s
 
 void testBranchingOnEnqueueKernel(queue_t default_queue, unsigned flags, ndrange_t ndrange) {
     // Ensure `enqueue_kernel` can be branched upon.
 
     if (enqueue_kernel(default_queue, flags, ndrange, ^(void) {}))
         (void)0;
     // CHECK: [[P:%[0-9]+]] = call spir_func i32 @__enqueue_kernel
     // CHECK-NEXT: [[Q:%[a-z0-9]+]] = icmp ne i32 [[P]], 0
void testBranchinOnPipeOperations(read_only pipe int r, write_only pipe int w, global int* ptr) {
     // CHECK-NEXT: icmp ne i32 [[R]], 0
 }
 
 void testBranchingOnAddressSpaceCast(generic long* ptr) {
     // Verify that pointer types are properly casted, respecting address spaces.
 
     if (to_global(ptr))
         (void)0;
-    // CHECK:       [[P:%[0-9]+]] = call spir_func [[GLOBAL_VOID:i8 addrspace\(1\)\*]] @__to_global([[GENERIC_VOID:i8 addrspace\(4\)\*]] {{%[0-9]+}})
-    // CHECK-NEXT:  [[Q:%[0-9]+]] = bitcast [[GLOBAL_VOID]] [[P]] to [[GLOBAL_i64:i64 addrspace\(1\)\*]]
-    // CHECK-NEXT:  [[BOOL:%[a-z0-9]+]] = icmp ne [[GLOBAL_i64]] [[Q]], null
+    // CHECK:       [[P:%[0-9]+]] = call spir_func [[GLOBAL_VOID:ptr addrspace\(1\)]] @__to_global([[GENERIC_VOID:ptr addrspace\(4\)]] {{%[0-9]+}})
+    // CHECK-NEXT:  [[BOOL:%[a-z0-9]+]] = icmp ne ptr addrspace(1) [[P]], null
     // CHECK-NEXT:  br i1 [[BOOL]]
 
     if (to_local(ptr))
         (void)0;
-    // CHECK:       [[P:%[0-9]+]] = call spir_func [[LOCAL_VOID:i8 addrspace\(3\)\*]] @__to_local([[GENERIC_VOID]] {{%[0-9]+}})
-    // CHECK-NEXT:  [[Q:%[0-9]+]] = bitcast [[LOCAL_VOID]] [[P]] to [[LOCAL_i64:i64 addrspace\(3\)\*]]
-    // CHECK-NEXT:  [[BOOL:%[a-z0-9]+]] = icmp ne [[LOCAL_i64]] [[Q]], null
+    // CHECK:       [[P:%[0-9]+]] = call spir_func [[LOCAL_VOID:ptr addrspace\(3\)]] @__to_local([[GENERIC_VOID]] {{%[0-9]+}})
+    // CHECK-NEXT:  [[BOOL:%[a-z0-9]+]] = icmp ne ptr addrspace(3) [[P]], null
     // CHECK-NEXT:  br i1 [[BOOL]]
 
     if (to_private(ptr))
         (void)0;
-    // CHECK:       [[P:%[0-9]+]] = call spir_func [[PRIVATE_VOID:i8\*]] @__to_private([[GENERIC_VOID]] {{%[0-9]+}})
-    // CHECK-NEXT:  [[Q:%[0-9]+]] = bitcast [[PRIVATE_VOID]] [[P]] to [[PRIVATE_i64:i64\*]]
-    // CHECK-NEXT:  [[BOOL:%[a-z0-9]+]] = icmp ne [[PRIVATE_i64]] [[Q]], null
+    // CHECK:       [[P:%[0-9]+]] = call spir_func [[PRIVATE_VOID:ptr]] @__to_private([[GENERIC_VOID]] {{%[0-9]+}})
+    // CHECK-NEXT:  [[BOOL:%[a-z0-9]+]] = icmp ne ptr [[P]], null
     // CHECK-NEXT:  br i1 [[BOOL]]
 }
 

clang/test/CodeGenOpenCL/cast_image.cl

-// RUN: %clang_cc1 -no-opaque-pointers -emit-llvm -o - -triple amdgcn--amdhsa %s | FileCheck --check-prefix=AMDGCN %s
-// RUN: %clang_cc1 -no-opaque-pointers -emit-llvm -o - -triple x86_64-unknown-unknown %s | FileCheck --check-prefix=X86 %s
+// RUN: %clang_cc1 -emit-llvm -o - -triple amdgcn--amdhsa %s | FileCheck --check-prefix=AMDGCN %s
+// RUN: %clang_cc1 -emit-llvm -o - -triple x86_64-unknown-unknown %s | FileCheck --check-prefix=X86 %s
 
 #ifdef __AMDGCN__
 
 constant int* convert(image2d_t img) {
-  // AMDGCN: bitcast %opencl.image2d_ro_t addrspace(4)* %img to i32 addrspace(4)*
+  // AMDGCN: ret ptr addrspace(4) %img
   return __builtin_astype(img, constant int*);
 }
 
 #else
 
 global int* convert(image2d_t img) {
-  // X86: bitcast %opencl.image2d_ro_t* %img to i32*
+  // X86: ret ptr %img
   return __builtin_astype(img, global int*);
 }
 
 #endif

clang/test/CodeGenOpenCL/const-str-array-decay.cl

-// RUN: %clang_cc1 -no-opaque-pointers %s -emit-llvm -o - -ffake-address-space-map | FileCheck %s
+// RUN: %clang_cc1 %s -emit-llvm -o - -ffake-address-space-map | FileCheck %s
 
 int test_func(constant char* foo);
 
 kernel void str_array_decy() {
   test_func("Test string literal");
 }
 
-// CHECK: i8 addrspace(2)* noundef getelementptr inbounds ([20 x i8], [20 x i8] addrspace(2)*
-// CHECK-NOT: addrspacecast
-
+// CHECK: call i32 @test_func(ptr addrspace(2) noundef @{{.*}})

clang/test/CodeGenOpenCL/kernels-have-spir-cc-by-default.cl

-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL1.2 -emit-llvm -triple x86_64-unknown-unknown -o - | FileCheck %s
-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL1.2 -emit-llvm -triple amdgcn-unknown-unknown -o - | FileCheck -check-prefixes=AMDGCN %s
+// RUN: %clang_cc1 %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 amdgcn-unknown-unknown -o - | FileCheck -check-prefixes=AMDGCN %s
 // Test that the kernels always use the SPIR calling convention
 // to have unambiguous mapping of arguments to feasibly implement
 // clSetKernelArg().
 
 typedef struct int_single {
     int a;
 } int_single;
 
typedef struct test_struct {
     float elementF;
     short elementG;
     double elementH;
 } test_struct;
 
 kernel void test_single(int_single input, global int* output) {
 // CHECK: spir_kernel
 // AMDGCN: define{{.*}} amdgpu_kernel void @test_single
-// CHECK: struct.int_single* nocapture {{.*}} byval(%struct.int_single)
-// CHECK: i32* nocapture noundef writeonly align 4 %output
+// CHECK: ptr nocapture {{.*}} byval(%struct.int_single)
+// CHECK: ptr nocapture noundef writeonly align 4 %output
  output[0] = input.a;
 }
 
 kernel void test_pair(int_pair input, global int* output) {
 // CHECK: spir_kernel
 // AMDGCN: define{{.*}} amdgpu_kernel void @test_pair
-// CHECK: struct.int_pair* nocapture {{.*}} byval(%struct.int_pair)
-// CHECK: i32* nocapture noundef writeonly align 4 %output
+// CHECK: ptr nocapture {{.*}} byval(%struct.int_pair)
+// CHECK: ptr nocapture noundef writeonly align 4 %output
  output[0] = (int)input.a;
  output[1] = (int)input.b;
 }
 
 kernel void test_kernel(test_struct input, global int* output) {
 // CHECK: spir_kernel
 // AMDGCN: define{{.*}} amdgpu_kernel void @test_kernel
-// CHECK: struct.test_struct* nocapture {{.*}} byval(%struct.test_struct)
-// CHECK: i32* nocapture noundef writeonly align 4 %output
+// CHECK: ptr nocapture {{.*}} byval(%struct.test_struct)
+// CHECK: ptr nocapture noundef writeonly align 4 %output
  output[0] = input.elementA;
  output[1] = input.elementB;
  output[2] = (int)input.elementC;
  output[3] = (int)input.elementD;
  output[4] = (int)input.elementE;
  output[5] = (int)input.elementF;
  output[6] = (int)input.elementG;
  output[7] = (int)input.elementH;
 };
 
 void test_function(int_pair input, global int* output) {
 // CHECK-NOT: spir_kernel
 // AMDGCN-NOT: define{{.*}} amdgpu_kernel void @test_function
-// CHECK: i64 %input.coerce0, i64 %input.coerce1, i32* nocapture noundef writeonly %output
+// CHECK: i64 %input.coerce0, i64 %input.coerce1, ptr nocapture noundef writeonly %output
  output[0] = (int)input.a;
  output[1] = (int)input.b;
 }

clang/test/CodeGenOpenCL/no-half.cl

-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=cl2.0 -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=cl1.2 -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=cl1.1 -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=cl2.0 -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=cl1.2 -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=cl1.1 -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
 
 #pragma OPENCL EXTENSION cl_khr_fp64:enable
 
-// CHECK-LABEL: @test_store_float(float noundef %foo, half addrspace({{.}}){{.*}} %bar)
+// CHECK-LABEL: @test_store_float(float noundef %foo, ptr addrspace({{.}}){{.*}} %bar)
 __kernel void test_store_float(float foo, __global half* bar)
 {
   __builtin_store_halff(foo, bar);
 // CHECK: [[HALF_VAL:%.*]] = fptrunc float %foo to half
-// CHECK: store half [[HALF_VAL]], half addrspace({{.}})* %bar, align 2
+// CHECK: store half [[HALF_VAL]], ptr addrspace({{.}}) %bar, align 2
 }
 
-// CHECK-LABEL: @test_store_double(double noundef %foo, half addrspace({{.}}){{.*}} %bar)
+// CHECK-LABEL: @test_store_double(double noundef %foo, ptr addrspace({{.}}){{.*}} %bar)
 __kernel void test_store_double(double foo, __global half* bar)
 {
   __builtin_store_half(foo, bar);
 // CHECK: [[HALF_VAL:%.*]] = fptrunc double %foo to half
-// CHECK: store half [[HALF_VAL]], half addrspace({{.}})* %bar, align 2
+// CHECK: store half [[HALF_VAL]], ptr addrspace({{.}}) %bar, align 2
 }
 
-// CHECK-LABEL: @test_load_float(float addrspace({{.}}){{.*}} %foo, half addrspace({{.}}){{.*}} %bar)
+// CHECK-LABEL: @test_load_float(ptr addrspace({{.}}){{.*}} %foo, ptr addrspace({{.}}){{.*}} %bar)
 __kernel void test_load_float(__global float* foo, __global half* bar)
 {
   foo[0] = __builtin_load_halff(bar);
-// CHECK: [[HALF_VAL:%.*]] = load half, half addrspace({{.}})* %bar
+// CHECK: [[HALF_VAL:%.*]] = load half, ptr addrspace({{.}}) %bar
 // CHECK: [[FULL_VAL:%.*]] = fpext half [[HALF_VAL]] to float
-// CHECK: store float [[FULL_VAL]], float addrspace({{.}})* %foo
+// CHECK: store float [[FULL_VAL]], ptr addrspace({{.}}) %foo
 }
 
-// CHECK-LABEL: @test_load_double(double addrspace({{.}}){{.*}} %foo, half addrspace({{.}}){{.*}} %bar)
+// CHECK-LABEL: @test_load_double(ptr addrspace({{.}}){{.*}} %foo, ptr addrspace({{.}}){{.*}} %bar)
 __kernel void test_load_double(__global double* foo, __global half* bar)
 {
   foo[0] = __builtin_load_half(bar);
-// CHECK: [[HALF_VAL:%.*]] = load half, half addrspace({{.}})* %bar
+// CHECK: [[HALF_VAL:%.*]] = load half, ptr addrspace({{.}}) %bar
 // CHECK: [[FULL_VAL:%.*]] = fpext half [[HALF_VAL]] to double
-// CHECK: store double [[FULL_VAL]], double addrspace({{.}})* %foo
+// CHECK: store double [[FULL_VAL]], ptr addrspace({{.}}) %foo
 }

clang/test/CodeGenOpenCL/pipe_builtin.cl

-// RUN: %clang_cc1 -no-opaque-pointers -triple %itanium_abi_triple -emit-llvm -cl-ext=+cl_khr_subgroups -O0 -cl-std=clc++ -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple %itanium_abi_triple -emit-llvm -cl-ext=+cl_khr_subgroups -O0 -cl-std=clc++ -o - %s | FileCheck %s
 // FIXME: Add MS ABI manglings of OpenCL things and remove %itanium_abi_triple
 // above to support OpenCL in the MS C++ ABI.
 
-// CHECK-DAG: %opencl.pipe_ro_t = type opaque
-// CHECK-DAG: %opencl.pipe_wo_t = type opaque
-// CHECK-DAG: %opencl.reserve_id_t = type opaque
-
 #pragma OPENCL EXTENSION cl_khr_subgroups : enable
 
 void test1(read_only pipe int p, global int *ptr) {
-  // CHECK: call i32 @__read_pipe_2(%opencl.pipe_ro_t* %{{.*}}, i8* %{{.*}}, i32 4, i32 4)
+  // CHECK: call i32 @__read_pipe_2(ptr %{{.*}}, ptr %{{.*}}, i32 4, i32 4)
   read_pipe(p, ptr);
-  // CHECK: call %opencl.reserve_id_t* @__reserve_read_pipe(%opencl.pipe_ro_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4)
+  // CHECK: call ptr @__reserve_read_pipe(ptr %{{.*}}, i32 {{.*}}, i32 4, i32 4)
   reserve_id_t rid = reserve_read_pipe(p, 2);
-  // CHECK: call i32 @__read_pipe_4(%opencl.pipe_ro_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 {{.*}}, i8* %{{.*}}, i32 4, i32 4)
+  // CHECK: call i32 @__read_pipe_4(ptr %{{.*}}, ptr %{{.*}}, i32 {{.*}}, ptr %{{.*}}, i32 4, i32 4)
   read_pipe(p, rid, 2, ptr);
-  // CHECK: call void @__commit_read_pipe(%opencl.pipe_ro_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4)
+  // CHECK: call void @__commit_read_pipe(ptr %{{.*}}, ptr %{{.*}}, i32 4, i32 4)
   commit_read_pipe(p, rid);
 }
 
 void test2(write_only pipe int p, global int *ptr) {
-  // CHECK: call i32 @__write_pipe_2(%opencl.pipe_wo_t* %{{.*}}, i8* %{{.*}}, i32 4, i32 4)
+  // CHECK: call i32 @__write_pipe_2(ptr %{{.*}}, ptr %{{.*}}, i32 4, i32 4)
   write_pipe(p, ptr);
-  // CHECK: call %opencl.reserve_id_t* @__reserve_write_pipe(%opencl.pipe_wo_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4)
+  // CHECK: call ptr @__reserve_write_pipe(ptr %{{.*}}, i32 {{.*}}, i32 4, i32 4)
   reserve_id_t rid = reserve_write_pipe(p, 2);
-  // CHECK: call i32 @__write_pipe_4(%opencl.pipe_wo_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 {{.*}}, i8* %{{.*}}, i32 4, i32 4)
+  // CHECK: call i32 @__write_pipe_4(ptr %{{.*}}, ptr %{{.*}}, i32 {{.*}}, ptr %{{.*}}, i32 4, i32 4)
   write_pipe(p, rid, 2, ptr);
-  // CHECK: call void @__commit_write_pipe(%opencl.pipe_wo_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4)
+  // CHECK: call void @__commit_write_pipe(ptr %{{.*}}, ptr %{{.*}}, i32 4, i32 4)
   commit_write_pipe(p, rid);
 }
 
 void test3(read_only pipe int p, global int *ptr) {
-  // CHECK: call %opencl.reserve_id_t* @__work_group_reserve_read_pipe(%opencl.pipe_ro_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4)
+  // CHECK: call ptr @__work_group_reserve_read_pipe(ptr %{{.*}}, i32 {{.*}}, i32 4, i32 4)
   reserve_id_t rid = work_group_reserve_read_pipe(p, 2);
-  // CHECK: call void @__work_group_commit_read_pipe(%opencl.pipe_ro_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4)
+  // CHECK: call void @__work_group_commit_read_pipe(ptr %{{.*}}, ptr %{{.*}}, i32 4, i32 4)
   work_group_commit_read_pipe(p, rid);
 }
 
 void test4(write_only pipe int p, global int *ptr) {
-  // CHECK: call %opencl.reserve_id_t* @__work_group_reserve_write_pipe(%opencl.pipe_wo_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4)
+  // CHECK: call ptr @__work_group_reserve_write_pipe(ptr %{{.*}}, i32 {{.*}}, i32 4, i32 4)
   reserve_id_t rid = work_group_reserve_write_pipe(p, 2);
-  // CHECK: call void @__work_group_commit_write_pipe(%opencl.pipe_wo_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4)
+  // CHECK: call void @__work_group_commit_write_pipe(ptr %{{.*}}, ptr %{{.*}}, i32 4, i32 4)
   work_group_commit_write_pipe(p, rid);
 }
 
 void test5(read_only pipe int p, global int *ptr) {
-  // CHECK: call %opencl.reserve_id_t* @__sub_group_reserve_read_pipe(%opencl.pipe_ro_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4)
+  // CHECK: call ptr @__sub_group_reserve_read_pipe(ptr %{{.*}}, i32 {{.*}}, i32 4, i32 4)
   reserve_id_t rid = sub_group_reserve_read_pipe(p, 2);
-  // CHECK: call void @__sub_group_commit_read_pipe(%opencl.pipe_ro_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4)
+  // CHECK: call void @__sub_group_commit_read_pipe(ptr %{{.*}}, ptr %{{.*}}, i32 4, i32 4)
   sub_group_commit_read_pipe(p, rid);
 }
 
 void test6(write_only pipe int p, global int *ptr) {
-  // CHECK: call %opencl.reserve_id_t* @__sub_group_reserve_write_pipe(%opencl.pipe_wo_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4)
+  // CHECK: call ptr @__sub_group_reserve_write_pipe(ptr %{{.*}}, i32 {{.*}}, i32 4, i32 4)
   reserve_id_t rid = sub_group_reserve_write_pipe(p, 2);
-  // CHECK: call void @__sub_group_commit_write_pipe(%opencl.pipe_wo_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4)
+  // CHECK: call void @__sub_group_commit_write_pipe(ptr %{{.*}}, ptr %{{.*}}, i32 4, i32 4)
   sub_group_commit_write_pipe(p, rid);
 }
 
 void test7(read_only pipe int p, global int *ptr) {
-  // CHECK: call i32 @__get_pipe_num_packets_ro(%opencl.pipe_ro_t* %{{.*}}, i32 4, i32 4)
+  // CHECK: call i32 @__get_pipe_num_packets_ro(ptr %{{.*}}, i32 4, i32 4)
   *ptr = get_pipe_num_packets(p);
-  // CHECK: call i32 @__get_pipe_max_packets_ro(%opencl.pipe_ro_t* %{{.*}}, i32 4, i32 4)
+  // CHECK: call i32 @__get_pipe_max_packets_ro(ptr %{{.*}}, i32 4, i32 4)
   *ptr = get_pipe_max_packets(p);
 }
 
 void test8(write_only pipe int p, global int *ptr) {
-  // CHECK: call i32 @__get_pipe_num_packets_wo(%opencl.pipe_wo_t* %{{.*}}, i32 4, i32 4)
+  // CHECK: call i32 @__get_pipe_num_packets_wo(ptr %{{.*}}, i32 4, i32 4)
   *ptr = get_pipe_num_packets(p);
-  // CHECK: call i32 @__get_pipe_max_packets_wo(%opencl.pipe_wo_t* %{{.*}}, i32 4, i32 4)
+  // CHECK: call i32 @__get_pipe_max_packets_wo(ptr %{{.*}}, i32 4, i32 4)
   *ptr = get_pipe_max_packets(p);
 }

clang/test/CodeGenOpenCL/pipe_types.cl

-// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-linux-gnu -emit-llvm -O0 -cl-std=CL2.0 -o - %s | FileCheck --check-prefixes=CHECK,CHECK-STRUCT %s
-// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-linux-gnu -emit-llvm -O0 -cl-std=CL3.0 -cl-ext=-all,+__opencl_c_pipes,+__opencl_c_generic_address_space,+__opencl_c_program_scope_global_variables -o - %s | FileCheck --check-prefixes=CHECK %s
-// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-linux-gnu -emit-llvm -O0 -cl-std=CL3.0 -cl-ext=-all,+__opencl_c_pipes,+__opencl_c_generic_address_space -o - %s | FileCheck --check-prefixes=CHECK %s
-// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-linux-gnu -emit-llvm -O0 -cl-std=clc++2021 -cl-ext=-all,+__opencl_c_pipes,+__opencl_c_generic_address_space,+__opencl_c_program_scope_global_variables -o - %s | FileCheck --check-prefixes=CHECK %s
-// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-linux-gnu -emit-llvm -O0 -cl-std=clc++2021 -cl-ext=-all,+__opencl_c_pipes,+__opencl_c_generic_address_space -o - %s | FileCheck --check-prefixes=CHECK %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -O0 -cl-std=CL2.0 -o - %s | FileCheck --check-prefixes=CHECK,CHECK-STRUCT %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -O0 -cl-std=CL3.0 -cl-ext=-all,+__opencl_c_pipes,+__opencl_c_generic_address_space,+__opencl_c_program_scope_global_variables -o - %s | FileCheck --check-prefixes=CHECK %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -O0 -cl-std=CL3.0 -cl-ext=-all,+__opencl_c_pipes,+__opencl_c_generic_address_space -o - %s | FileCheck --check-prefixes=CHECK %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -O0 -cl-std=clc++2021 -cl-ext=-all,+__opencl_c_pipes,+__opencl_c_generic_address_space,+__opencl_c_program_scope_global_variables -o - %s | FileCheck --check-prefixes=CHECK %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -O0 -cl-std=clc++2021 -cl-ext=-all,+__opencl_c_pipes,+__opencl_c_generic_address_space -o - %s | FileCheck --check-prefixes=CHECK %s
 
-// CHECK: %opencl.pipe_ro_t = type opaque
-// CHECK: %opencl.pipe_wo_t = type opaque
 typedef unsigned char __attribute__((ext_vector_type(3))) uchar3;
 typedef int __attribute__((ext_vector_type(4))) int4;
 
 void test1(read_only pipe int p) {
-// CHECK: define{{.*}} void @{{.*}}test1{{.*}}(%opencl.pipe_ro_t* %p)
+// CHECK: define{{.*}} void @{{.*}}test1{{.*}}(ptr %p)
   reserve_id_t rid;
-// CHECK: %rid = alloca %opencl.reserve_id_t
+// CHECK: %rid = alloca ptr
 }
 
 void test2(write_only pipe float p) {
-// CHECK: define{{.*}} void @{{.*}}test2{{.*}}(%opencl.pipe_wo_t* %p)
+// CHECK: define{{.*}} void @{{.*}}test2{{.*}}(ptr %p)
 }
 
 void test3(read_only pipe const int p) {
-// CHECK: define{{.*}} void @{{.*}}test3{{.*}}(%opencl.pipe_ro_t* %p)
+// CHECK: define{{.*}} void @{{.*}}test3{{.*}}(ptr %p)
 }
 
 void test4(read_only pipe uchar3 p) {
-// CHECK: define{{.*}} void @{{.*}}test4{{.*}}(%opencl.pipe_ro_t* %p)
+// CHECK: define{{.*}} void @{{.*}}test4{{.*}}(ptr %p)
 }
 
 void test5(read_only pipe int4 p) {
-// CHECK: define{{.*}} void @{{.*}}test5{{.*}}(%opencl.pipe_ro_t* %p)
+// CHECK: define{{.*}} void @{{.*}}test5{{.*}}(ptr %p)
 }
 
 typedef read_only pipe int MyPipe;
 kernel void test6(MyPipe p) {
-// CHECK: define{{.*}} spir_kernel void @test6(%opencl.pipe_ro_t* %p)
+// CHECK: define{{.*}} spir_kernel void @test6(ptr %p)
 }
 
 struct Person {
   const char *Name;
   bool isFemale;
   int ID;
 };
 
 void test_reserved_read_pipe(global struct Person *SDst,
                              read_only pipe struct Person SPipe) {
   // CHECK-STRUCT: define{{.*}} void @test_reserved_read_pipe
   read_pipe (SPipe, SDst);
-  // CHECK-STRUCT: call i32 @__read_pipe_2(%opencl.pipe_ro_t* %{{.*}}, i8* %{{.*}}, i32 16, i32 8)
+  // CHECK-STRUCT: call i32 @__read_pipe_2(ptr %{{.*}}, ptr %{{.*}}, i32 16, i32 8)
   read_pipe (SPipe, SDst);
-  // CHECK-STRUCT: call i32 @__read_pipe_2(%opencl.pipe_ro_t* %{{.*}}, i8* %{{.*}}, i32 16, i32 8)
+  // CHECK-STRUCT: call i32 @__read_pipe_2(ptr %{{.*}}, ptr %{{.*}}, i32 16, i32 8)
 }

clang/test/CodeGenOpenCL/to_addr_builtin.cl

-// RUN: %clang_cc1 -no-opaque-pointers -triple spir-unknown-unknown -emit-llvm -O0 -cl-std=clc++ -o - %s | FileCheck %s
-// RUN: %clang_cc1 -no-opaque-pointers -triple spir-unknown-unknown -emit-llvm -O0 -cl-std=cl2.0 -o - %s | FileCheck %s
-// RUN: %clang_cc1 -no-opaque-pointers -triple spir-unknown-unknown -emit-llvm -O0 -cl-std=cl3.0 -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple spir-unknown-unknown -emit-llvm -O0 -cl-std=clc++ -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple spir-unknown-unknown -emit-llvm -O0 -cl-std=cl2.0 -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple spir-unknown-unknown -emit-llvm -O0 -cl-std=cl3.0 -o - %s | FileCheck %s
 
-// CHECK: %[[A:.*]] = type { float, float, float }
 typedef struct {
   float x,y,z;
 } A;
 typedef private A *PA;
 typedef global A *GA;
 
 void test(void) {
   global int *glob;
   local int *loc;
   private int *priv;
   generic int *gen;
 
-  //CHECK: %[[ARG:.*]] = addrspacecast i32 addrspace(1)* %{{.*}} to i8 addrspace(4)*
-  //CHECK: %[[RET:.*]] = call spir_func i8 addrspace(1)* @__to_global(i8 addrspace(4)* %[[ARG]])
-  //CHECK: %{{.*}} = bitcast i8 addrspace(1)* %[[RET]] to i32 addrspace(1)*
+  //CHECK: %[[ARG:.*]] = addrspacecast ptr addrspace(1) %{{.*}} to ptr addrspace(4)
+  //CHECK: %[[RET:.*]] = call spir_func ptr addrspace(1) @__to_global(ptr addrspace(4) %[[ARG]])
+  //CHECK: store ptr addrspace(1) %[[RET]], ptr %glob
   glob = to_global(glob);
   
-  //CHECK: %[[ARG:.*]] = addrspacecast i32 addrspace(3)* %{{.*}} to i8 addrspace(4)*
-  //CHECK: %[[RET:.*]] = call spir_func i8 addrspace(1)* @__to_global(i8 addrspace(4)* %[[ARG]])
-  //CHECK: %{{.*}} = bitcast i8 addrspace(1)* %[[RET]] to i32 addrspace(1)*
+  //CHECK: %[[ARG:.*]] = addrspacecast ptr addrspace(3) %{{.*}} to ptr addrspace(4)
+  //CHECK: %[[RET:.*]] = call spir_func ptr addrspace(1) @__to_global(ptr addrspace(4) %[[ARG]])
+  //CHECK: store ptr addrspace(1) %[[RET]], ptr %glob
   glob = to_global(loc);
  
-  //CHECK: %[[ARG:.*]] = addrspacecast i32* %{{.*}} to i8 addrspace(4)*
-  //CHECK: %[[RET:.*]] = call spir_func i8 addrspace(1)* @__to_global(i8 addrspace(4)* %[[ARG]])
-  //CHECK: %{{.*}} = bitcast i8 addrspace(1)* %[[RET]] to i32 addrspace(1)*
+  //CHECK: %[[ARG:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(4)
+  //CHECK: %[[RET:.*]] = call spir_func ptr addrspace(1) @__to_global(ptr addrspace(4) %[[ARG]])
+  //CHECK: store ptr addrspace(1) %[[RET]], ptr %glob
   glob = to_global(priv);
  
-  //CHECK: %[[ARG:.*]] = bitcast i32 addrspace(4)* %{{.*}} to i8 addrspace(4)*
-  //CHECK: %[[RET:.*]] = call spir_func i8 addrspace(1)* @__to_global(i8 addrspace(4)* %[[ARG]])
-  //CHECK: %{{.*}} = bitcast i8 addrspace(1)* %[[RET]] to i32 addrspace(1)*
+  //CHECK: %[[ARG:.*]] = load ptr addrspace(4), ptr %gen
+  //CHECK: %[[RET:.*]] = call spir_func ptr addrspace(1) @__to_global(ptr addrspace(4) %[[ARG]])
+  //CHECK: store ptr addrspace(1) %[[RET]], ptr %glob
   glob = to_global(gen);
   
-  //CHECK: %[[ARG:.*]] = addrspacecast i32 addrspace(1)* %{{.*}} to i8 addrspace(4)*
-  //CHECK: %[[RET:.*]] = call spir_func i8 addrspace(3)* @__to_local(i8 addrspace(4)* %[[ARG]])
-  //CHECK: %{{.*}} = bitcast i8 addrspace(3)* %[[RET]] to i32 addrspace(3)*
+  //CHECK: %[[ARG:.*]] = addrspacecast ptr addrspace(1) %{{.*}} to ptr addrspace(4)
+  //CHECK: %[[RET:.*]] = call spir_func ptr addrspace(3) @__to_local(ptr addrspace(4) %[[ARG]])
+  //CHECK: store ptr addrspace(3) %[[RET]], ptr %loc
   loc = to_local(glob);
 
-  //CHECK: %[[ARG:.*]] = addrspacecast i32 addrspace(3)* %{{.*}} to i8 addrspace(4)*
-  //CHECK: %[[RET:.*]] = call spir_func i8 addrspace(3)* @__to_local(i8 addrspace(4)* %[[ARG]])
-  //CHECK: %{{.*}} = bitcast i8 addrspace(3)* %[[RET]] to i32 addrspace(3)*
+  //CHECK: %[[ARG:.*]] = addrspacecast ptr addrspace(3) %{{.*}} to ptr addrspace(4)
+  //CHECK: %[[RET:.*]] = call spir_func ptr addrspace(3) @__to_local(ptr addrspace(4) %[[ARG]])
+  //CHECK: store ptr addrspace(3) %[[RET]], ptr %loc
   loc = to_local(loc);
 
-  //CHECK: %[[ARG:.*]] = addrspacecast i32* %{{.*}} to i8 addrspace(4)*
-  //CHECK: %[[RET:.*]] = call spir_func i8 addrspace(3)* @__to_local(i8 addrspace(4)* %[[ARG]])
-  //CHECK: %{{.*}} = bitcast i8 addrspace(3)* %[[RET]] to i32 addrspace(3)*
+  //CHECK: %[[ARG:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(4)
+  //CHECK: %[[RET:.*]] = call spir_func ptr addrspace(3) @__to_local(ptr addrspace(4) %[[ARG]])
+  //CHECK: store ptr addrspace(3) %[[RET]], ptr %loc
   loc = to_local(priv);
 
-  //CHECK: %[[ARG:.*]] = bitcast i32 addrspace(4)* %{{.*}} to i8 addrspace(4)*
-  //CHECK: %[[RET:.*]] = call spir_func i8 addrspace(3)* @__to_local(i8 addrspace(4)* %[[ARG]])
-  //CHECK: %{{.*}} = bitcast i8 addrspace(3)* %[[RET]] to i32 addrspace(3)*
+  //CHECK: %[[ARG:.*]] = load ptr addrspace(4), ptr %gen
+  //CHECK: %[[RET:.*]] = call spir_func ptr addrspace(3) @__to_local(ptr addrspace(4) %[[ARG]])
+  //CHECK: store ptr addrspace(3) %[[RET]], ptr %loc
   loc = to_local(gen);
 
-  //CHECK: %[[ARG:.*]] = addrspacecast i32 addrspace(1)* %{{.*}} to i8 addrspace(4)*
-  //CHECK: %[[RET:.*]] = call spir_func i8* @__to_private(i8 addrspace(4)* %[[ARG]])
-  //CHECK: %{{.*}} = bitcast i8* %[[RET]] to i32*
+  //CHECK: %[[ARG:.*]] = addrspacecast ptr addrspace(1) %{{.*}} to ptr addrspace(4)
+  //CHECK: %[[RET:.*]] = call spir_func ptr @__to_private(ptr addrspace(4) %[[ARG]])
+  //CHECK: store ptr %[[RET]], ptr %priv
   priv = to_private(glob);
 
-  //CHECK: %[[ARG:.*]] = addrspacecast i32 addrspace(3)* %{{.*}} to i8 addrspace(4)*
-  //CHECK: %[[RET:.*]] = call spir_func i8* @__to_private(i8 addrspace(4)* %[[ARG]])
-  //CHECK: %{{.*}} = bitcast i8* %[[RET]] to i32*
+  //CHECK: %[[ARG:.*]] = addrspacecast ptr addrspace(3) %{{.*}} to ptr addrspace(4)
+  //CHECK: %[[RET:.*]] = call spir_func ptr @__to_private(ptr addrspace(4) %[[ARG]])
+  //CHECK: store ptr %[[RET]], ptr %priv
   priv = to_private(loc);
 
-  //CHECK: %[[ARG:.*]] = addrspacecast i32* %{{.*}} to i8 addrspace(4)*
-  //CHECK: %[[RET:.*]] = call spir_func i8* @__to_private(i8 addrspace(4)* %[[ARG]])
-  //CHECK: %{{.*}} = bitcast i8* %[[RET]] to i32*
+  //CHECK: %[[ARG:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(4)
+  //CHECK: %[[RET:.*]] = call spir_func ptr @__to_private(ptr addrspace(4) %[[ARG]])
+  //CHECK: store ptr %[[RET]], ptr %priv
   priv = to_private(priv);
 
-  //CHECK: %[[ARG:.*]] = bitcast i32 addrspace(4)* %{{.*}} to i8 addrspace(4)*
-  //CHECK: %[[RET:.*]] = call spir_func i8* @__to_private(i8 addrspace(4)* %[[ARG]])
-  //CHECK: %{{.*}} = bitcast i8* %[[RET]] to i32*
+  //CHECK: %[[ARG:.*]] = load ptr addrspace(4), ptr %gen
+  //CHECK: %[[RET:.*]] = call spir_func ptr @__to_private(ptr addrspace(4) %[[ARG]])
+  //CHECK: store ptr %[[RET]], ptr %priv
   priv = to_private(gen);
 
-  //CHECK: %[[ARG:.*]] = addrspacecast %[[A]]* %{{.*}} to i8 addrspace(4)*
-  //CHECK: %[[RET:.*]] = call spir_func i8 addrspace(1)* @__to_global(i8 addrspace(4)* %[[ARG]])
-  //CHECK: %{{.*}} = bitcast i8 addrspace(1)* %[[RET]] to %[[A]] addrspace(1)*
+  //CHECK: %[[ARG:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(4)
+  //CHECK: %[[RET:.*]] = call spir_func ptr addrspace(1) @__to_global(ptr addrspace(4) %[[ARG]])
+  //CHECK: store ptr addrspace(1) %[[RET]], ptr %gA
   PA pA;
   GA gA = to_global(pA);
 
   //CHECK-NOT: addrspacecast
   //CHECK-NOT: bitcast
-  //CHECK: call spir_func i8 addrspace(1)* @__to_global(i8 addrspace(4)* %{{.*}})
+  //CHECK: call spir_func ptr addrspace(1) @__to_global(ptr addrspace(4) %{{.*}})
   //CHECK-NOT: addrspacecast
   //CHECK-NOT: bitcast
   generic void *gen_v;
   global void *glob_v = to_global(gen_v);
 }

clang/test/CodeGenSYCL/field-annotate-addr-space.cpp

-// RUN: %clang_cc1 -no-opaque-pointers -triple spir64 -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
 
 // CHECK: [[ANNOT:.+]] = private unnamed_addr addrspace(1) constant {{.*}}c"my_annotation\00"
 
 struct HasField {
   // This caused an assertion on creating a bitcast here,
   // since the address space didn't match.
   [[clang::annotate("my_annotation")]]
   int *a;
 };
 
 void foo(int *b) {
   struct HasField f;
-  // CHECK: %[[A:.+]] = getelementptr inbounds %struct.HasField, %struct.HasField addrspace(4)* %{{.+}}
-  // CHECK: %[[BITCAST:.+]] = bitcast i32 addrspace(4)* addrspace(4)* %[[A]] to i8 addrspace(4)*
-  // CHECK: %[[CALL:.+]] = call i8 addrspace(4)* @llvm.ptr.annotation.p4i8.p1i8(i8 addrspace(4)* %[[BITCAST]], i8 addrspace(1)* getelementptr inbounds ([14 x i8], [14 x i8] addrspace(1)* [[ANNOT]]
-  // CHECK: bitcast i8 addrspace(4)* %[[CALL]] to i32 addrspace(4)* addrspace(4)*
+  // CHECK: %[[A:.+]] = getelementptr inbounds %struct.HasField, ptr addrspace(4) %{{.+}}
+  // CHECK: %[[CALL:.+]] = call ptr addrspace(4) @llvm.ptr.annotation.p4.p1(ptr addrspace(4) %[[A]], ptr addrspace(1) [[ANNOT]]
+  // CHECK: store ptr addrspace(4) %{{[0-9]+}}, ptr addrspace(4) %[[CALL]]
   f.a = b;
 }

clang/test/CodeGenSYCL/unique_stable_name.cpp

-// RUN: %clang_cc1 -no-opaque-pointers -triple spir64-unknown-unknown-sycldevice -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
 // CHECK: @[[LAMBDA_KERNEL3:[^\w]+]] = private unnamed_addr addrspace(1) constant [[LAMBDA_K3_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ4mainEUlPZ4mainEUlvE_E_\00"
 // CHECK: @[[INT1:[^\w]+]] = private unnamed_addr addrspace(1) constant [[INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSi\00"
 // CHECK: @[[STRING:[^\w]+]] = private unnamed_addr addrspace(1) constant [[STRING_SIZE:\[[0-9]+ x i8\]]] c"_ZTSAppL_ZZ4mainE1jE_i\00",
 // CHECK: @[[INT2:[^\w]+]] = private unnamed_addr addrspace(1) constant [[INT_SIZE]] c"_ZTSi\00"
 // CHECK: @[[LAMBDA_X:[^\w]+]] = private unnamed_addr addrspace(1) constant [[LAMBDA_X_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE_\00"
 // CHECK: @[[MACRO_X:[^\w]+]] = private unnamed_addr addrspace(1) constant [[MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE0_\00"
 // CHECK: @[[MACRO_Y:[^\w]+]] =  private unnamed_addr addrspace(1) constant [[MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE1_\00"
 // CHECK: @{{.*}} = private unnamed_addr addrspace(1) constant [32 x i8] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE2_\00", align 1

 
 template <typename KernelName, typename KernelType>
 [[clang::sycl_kernel]] void kernel_single_task(KernelType kernelFunc) {
   kernelFunc();
 }
 
 int main() {
   kernel_single_task<class kernel2>(func<Derp>);
-  // CHECK: call spir_func void @_Z18kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_(i8 addrspace(4)* ()* noundef @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv)
+  // CHECK: call spir_func void @_Z18kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_(ptr noundef @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv)
 
   auto l1 = []() { return 1; };
   auto l2 = [](decltype(l1) *l = nullptr) { return 2; };
   kernel_single_task<class kernel3>(l2);
   puts(__builtin_sycl_unique_stable_name(decltype(l2)));
   // CHECK: call spir_func void @_Z18kernel_single_taskIZ4mainE7kernel3Z4mainEUlPZ4mainEUlvE_E_EvT0_
-  // CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[LAMBDA_K3_SIZE]], [[LAMBDA_K3_SIZE]] addrspace(1)* @[[LAMBDA_KERNEL3]], i32 0, i32 0) to i8 addrspace(4)*))
+  // CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[LAMBDA_KERNEL3]] to ptr addrspace(4)))
 
   constexpr const char str[] = "lalala";
   static_assert(__builtin_strcmp(__builtin_sycl_unique_stable_name(decltype(str)), "_ZTSA7_Kc\0") == 0, "unexpected mangling");
 
   int i = 0;
   puts(__builtin_sycl_unique_stable_name(decltype(i++)));
-  // CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]] addrspace(1)* @[[INT1]], i32 0, i32 0) to i8 addrspace(4)*))
+  // CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[INT1]] to ptr addrspace(4)))
 
   // FIXME: Ensure that j is incremented because VLAs are terrible.
   int j = 55;
   puts(__builtin_sycl_unique_stable_name(int[++j]));
-  // CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[STRING_SIZE]], [[STRING_SIZE]] addrspace(1)* @[[STRING]], i32 0, i32 0) to i8 addrspace(4)*))
+  // CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[STRING]] to ptr addrspace(4)))
 
   // CHECK: define internal spir_func void @_Z18kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_
-  // CHECK: declare spir_func noundef i8 addrspace(4)* @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv
+  // CHECK: declare spir_func noundef ptr addrspace(4) @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv
   // CHECK: define internal spir_func void @_Z18kernel_single_taskIZ4mainE7kernel3Z4mainEUlPZ4mainEUlvE_E_EvT0_
   // CHECK: define internal spir_func void @_Z18kernel_single_taskIZ4mainE6kernelZ4mainEUlvE0_EvT0_
 
   kernel_single_task<class kernel>(
       []() {
         puts(__builtin_sycl_unique_stable_name(int));
-        // CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]] addrspace(1)* @[[INT2]], i32 0, i32 0) to i8 addrspace(4)*))
+        // CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[INT2]] to ptr addrspace(4)))
 
         auto x = []() {};
         puts(__builtin_sycl_unique_stable_name(decltype(x)));
-        // CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]] addrspace(1)* @[[LAMBDA_X]], i32 0, i32 0) to i8 addrspace(4)*))
+        // CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[LAMBDA_X]] to ptr addrspace(4)))
 
         DEF_IN_MACRO();
-        // CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]] addrspace(1)* @[[MACRO_X]], i32 0, i32 0) to i8 addrspace(4)*))
-        // CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]] addrspace(1)* @[[MACRO_Y]], i32 0, i32 0) to i8 addrspace(4)*))
+        // CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[MACRO_X]] to ptr addrspace(4)))
+        // CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[MACRO_Y]] to ptr addrspace(4)))
 
         MACRO_CALLS_MACRO();
-        // CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]] addrspace(1)* @[[MACRO_MACRO_X]], i32 0, i32 0) to i8 addrspace(4)*))
-        // CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]] addrspace(1)* @[[MACRO_MACRO_Y]], i32 0, i32 0) to i8 addrspace(4)*))
+        // CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[MACRO_MACRO_X]] to ptr addrspace(4)))
+        // CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[MACRO_MACRO_Y]] to ptr addrspace(4)))
 
         template_param<int>();
         // CHECK: call spir_func void @_Z14template_paramIiEvv
 
         template_param<decltype(x)>();
         // CHECK: call spir_func void @_Z14template_paramIZZ4mainENKUlvE0_clEvEUlvE_Evv
 
         lambda_in_dependent_function<int>();
kernel_single_task<class kernel>(
         // CHECK: call spir_func void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv
 
         lambda_two_dep<decltype(z), decltype(y)>();
         // CHECK: call spir_func void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv
       });
 }
 
 // CHECK: define linkonce_odr spir_func void @_Z14template_paramIiEvv
-// CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]] addrspace(1)* @[[INT3]], i32 0, i32 0) to i8 addrspace(4)*))
+// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[INT3]] to ptr addrspace(4)))
 
 // CHECK: define internal spir_func void @_Z14template_paramIZZ4mainENKUlvE0_clEvEUlvE_Evv
-// CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[LAMBDA_SIZE]], [[LAMBDA_SIZE]] addrspace(1)* @[[LAMBDA]], i32 0, i32 0) to i8 addrspace(4)*))
+// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[LAMBDA]] to ptr addrspace(4)))
 
 // CHECK: define linkonce_odr spir_func void @_Z28lambda_in_dependent_functionIiEvv
-// CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[DEP_INT_SIZE]], [[DEP_INT_SIZE]] addrspace(1)* @[[LAMBDA_IN_DEP_INT]], i32 0, i32 0) to i8 addrspace(4)*))
+// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[LAMBDA_IN_DEP_INT]] to ptr addrspace(4)))
 
 // CHECK: define internal spir_func void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_Evv
-// CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[DEP_LAMBDA_SIZE]], [[DEP_LAMBDA_SIZE]] addrspace(1)* @[[LAMBDA_IN_DEP_X]], i32 0, i32 0) to i8 addrspace(4)*))
+// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[LAMBDA_IN_DEP_X]] to ptr addrspace(4)))
 
 // CHECK: define linkonce_odr spir_func void @_Z13lambda_no_depIidEvT_T0_(i32 noundef %a, double noundef %b)
-// CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[NO_DEP_LAMBDA_SIZE]], [[NO_DEP_LAMBDA_SIZE]] addrspace(1)* @[[LAMBDA_NO_DEP]], i32 0, i32 0) to i8 addrspace(4)*))
+// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[LAMBDA_NO_DEP]] to ptr addrspace(4)))
 
 // CHECK: define internal spir_func void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv
-// CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[DEP_LAMBDA1_SIZE]], [[DEP_LAMBDA1_SIZE]] addrspace(1)* @[[LAMBDA_TWO_DEP]], i32 0, i32 0) to i8 addrspace(4)*))
+// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[LAMBDA_TWO_DEP]] to ptr addrspace(4)))
 
 // CHECK: define internal spir_func void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv
-// CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[DEP_LAMBDA2_SIZE]], [[DEP_LAMBDA2_SIZE]] addrspace(1)* @[[LAMBDA_TWO_DEP2]], i32 0, i32 0) to i8 addrspace(4)*))
+// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[LAMBDA_TWO_DEP2]] to ptr addrspace(4)))

clang/test/PCH/arc-blocks.mm

-// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-apple-darwin10 -fobjc-arc -fblocks -std=c++1y -emit-pch %s -o %t
-// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-apple-darwin10 -fobjc-arc -fblocks -std=c++1y -include-pch %t -fobjc-avoid-heapify-local-blocks -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-apple-darwin10 -fobjc-arc -fblocks -std=c++1y -emit-pch %s -o %t
+// RUN: %clang_cc1 -triple x86_64-apple-darwin10 -fobjc-arc -fblocks -std=c++1y -include-pch %t -fobjc-avoid-heapify-local-blocks -emit-llvm -o - %s | FileCheck %s
 
 #ifndef HEADER_INCLUDED
 #define HEADER_INCLUDED
 
 namespace test_block_retain {
   typedef void (^BlockTy)();
   void foo1(id);
 
if (c)
       // @llvm.objc.retainBlock is called since 'b0' is declared in the outer scope.
       b0 = ^{ foo1(a); };
     b0();
   }
 }
 
 #else
 
-// CHECK: %[[STRUCT_BLOCK_DESCRIPTOR:.*]] = type { i64, i64 }
-
 namespace test_block_retain {
 // CHECK-LABEL: define linkonce_odr void @_ZN17test_block_retain14initializationEP11objc_object(
 // CHECK-NOT: call i8* @llvm.objc.retainBlock(
 
   void test_initialization(id a) {
     initialization(a);
   }
 
 // CHECK-LABEL: define{{.*}} void @_ZN17test_block_retain26test_assignmentConditionalEP11objc_objectb(
-// CHECK: %[[BLOCK:.*]] = alloca <{ i8*, i32, i32, i8*, %[[STRUCT_BLOCK_DESCRIPTOR]]*, i8* }>, align 8
-// CHECK: %[[V4:.*]] = bitcast <{ i8*, i32, i32, i8*, %[[STRUCT_BLOCK_DESCRIPTOR]]*, i8* }>* %[[BLOCK]] to void ()*
-// CHECK: %[[V5:.*]] = bitcast void ()* %[[V4]] to i8*
-// CHECK: call i8* @llvm.objc.retainBlock(i8* %[[V5]])
+// CHECK: %[[BLOCK:.*]] = alloca <{ ptr, i32, i32, ptr, ptr, ptr }>, align 8
+// CHECK: call ptr @llvm.objc.retainBlock(ptr %[[BLOCK]])
 
   void test_assignmentConditional(id a, bool c) {
     assignmentConditional(a, c);
   }
 }
 
 #endif

clang/test/Profile/objc-general.m

 // Test instrumentation of general constructs in objective C.
 
-// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-apple-macosx10.9 -main-file-name objc-general.m %s -o - -emit-llvm -fblocks -fprofile-instrument=clang | FileCheck -check-prefix=PGOGEN %s
+// RUN: %clang_cc1 -triple x86_64-apple-macosx10.9 -main-file-name objc-general.m %s -o - -emit-llvm -fblocks -fprofile-instrument=clang | FileCheck -check-prefix=PGOGEN %s
 
 // RUN: llvm-profdata merge %S/Inputs/objc-general.proftext -o %t.profdata
-// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-apple-macosx10.9 -main-file-name objc-general.m %s -o - -emit-llvm -fblocks -fprofile-instrument-use-path=%t.profdata 2>&1 | FileCheck -check-prefix=PGOUSE %s
+// RUN: %clang_cc1 -triple x86_64-apple-macosx10.9 -main-file-name objc-general.m %s -o - -emit-llvm -fblocks -fprofile-instrument-use-path=%t.profdata 2>&1 | FileCheck -check-prefix=PGOUSE %s
 
 // PGOUSE-NOT: warning: profile data may be out of date
 
 #ifdef HAVE_FOUNDATION
 
 // Use this to build an instrumented version to regenerate the input file.
 #import <Foundation/Foundation.h>
 

 
 @interface A : NSObject
 + (void)foreach: (NSArray *)array;
 @end
 
 @implementation A
 // PGOGEN: define {{.*}}+[A foreach:]
 // PGOUSE: define {{.*}}+[A foreach:]
-// PGOGEN: store {{.*}} @[[FRC]], i32 0, i32 0
+// PGOGEN: store {{.*}} @[[FRC]]
 + (void)foreach: (NSArray *)array
 {
   __block id result;
   // PGOGEN: store {{.*}} @[[FRC]], i32 0, i32 1
   // PGOUSE: br {{.*}} !prof ![[FR1:[0-9]+]]
   // PGOUSE: br {{.*}} !prof ![[FR2:[0-9]+]]
   for (id x in array) {
     // PGOGEN: define {{.*}}_block_invoke
     // PGOUSE: define {{.*}}_block_invoke
-    // PGOGEN: store {{.*}} @[[BLC]], i32 0, i32 0
+    // PGOGEN: store {{.*}} @[[BLC]]
     ^{
       static int init = 0;
       // PGOGEN: store {{.*}} @[[BLC]], i32 0, i32 1
       // PGOUSE: br {{.*}} !prof ![[BL1:[0-9]+]]
       if (init)
         result = x;
       init = 1;
     }();

clang/test/Rewriter/rewrite-super-message.mm

-// RUN: %clang_cc1 -no-opaque-pointers -x objective-c++ -Wno-return-type -fblocks -fms-extensions -rewrite-objc -fobjc-runtime=macosx-fragile-10.5 %s -o %t-rw.cpp
-// RUN: %clang_cc1 -no-opaque-pointers -fsyntax-only -Wno-address-of-temporary -DKEEP_ATTRIBUTES -D"id=struct objc_object *" -D"Class=struct objc_class *" -D"SEL=void*" -D"__declspec(X)=" -emit-llvm -o - %t-rw.cpp | FileCheck %t-rw.cpp
+// RUN: %clang_cc1 -x objective-c++ -Wno-return-type -fblocks -fms-extensions -rewrite-objc -fobjc-runtime=macosx-fragile-10.5 %s -o %t-rw.cpp
+// RUN: %clang_cc1 -fsyntax-only -Wno-address-of-temporary -DKEEP_ATTRIBUTES -D"id=struct objc_object *" -D"Class=struct objc_class *" -D"SEL=void*" -D"__declspec(X)=" -emit-llvm -o - %t-rw.cpp | FileCheck %t-rw.cpp
 // radar 7738453
 
 void *sel_registerName(const char *);
 
 @interface __NSCFType
 @end
 
 @interface __NSCFString : __NSCFType
 - (const char *)UTF8String;
 @end
 
 @implementation __NSCFString
 - (const char *)UTF8String {
     return (const char *)[super UTF8String];
 }
 @end
 
-// CHECK: call %struct.objc_class* @class_getSuperclass
+// CHECK: call ptr @class_getSuperclass
 
 @class NSZone;
 
 @interface NSObject {
 }
 
 + (id)allocWithZone:(NSZone *)zone;
 @end

clang/test/SemaObjC/arc-objc-lifetime-conflict.m

-// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-apple-darwin11 -fobjc-arc -fobjc-runtime-has-weak %s -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-apple-darwin11 -fobjc-arc -fobjc-runtime-has-weak %s -emit-llvm -o - | FileCheck %s
 
-// CHECK: bitcast {{.*}} %self_weak_s_w_s
-// CHECK-NEXT: llvm.objc.destroyWeak
-// CHECK-NEXT: bitcast {{.*}} %self_strong_w_s
-// CHECK-NEXT: llvm.objc.storeStrong
-// CHECK-NEXT: bitcast {{.*}} %self_weak_s
-// CHECK-NEXT: llvm.objc.destroyWeak
-// CHECK-NEXT: bitcast {{.*}} %self_weak_s3
-// CHECK-NEXT: llvm.objc.destroyWeak
-// CHECK-NEXT: bitcast {{.*}} %self_strong3
-// CHECK-NEXT: llvm.objc.storeStrong
-// CHECK-NEXT: bitcast {{.*}} %self_strong2
-// CHECK-NEXT: llvm.objc.storeStrong
-// CHECK-NEXT: bitcast {{.*}} %self_strong
-// CHECK-NEXT: llvm.objc.storeStrong
+// CHECK-DAG: llvm.objc.destroyWeak(ptr %self_weak_s_w_s)
+// CHECK-DAG: llvm.objc.storeStrong(ptr %self_strong_w_s, ptr null)
+// CHECK-DAG: llvm.objc.destroyWeak(ptr %self_weak_s)
+// CHECK-DAG: llvm.objc.destroyWeak(ptr %self_weak_s3)
+// CHECK-DAG: llvm.objc.storeStrong(ptr %self_strong3, ptr null)
+// CHECK-DAG: llvm.objc.storeStrong(ptr %self_strong2, ptr null)
+// CHECK-DAG: llvm.objc.storeStrong(ptr %self_strong, ptr null)
 @interface NSObject
 @end
 @interface A : NSObject
 @end
 @implementation A
 - (void)test {
   __attribute__((objc_ownership(strong))) __typeof__(self) self_strong;
   __attribute__((objc_ownership(strong))) __typeof__(self_strong) self_strong2;
   __attribute__((objc_ownership(strong))) __typeof__(self_strong2) self_strong3;
   __attribute__((objc_ownership(weak))) __typeof__(self_strong3) self_weak_s3;
  
   __attribute__((objc_ownership(weak))) __typeof__(self_strong) self_weak_s;
   __attribute__((objc_ownership(strong))) __typeof__(self_weak_s) self_strong_w_s;
   __attribute__((objc_ownership(weak))) __typeof__(self_strong_w_s) self_weak_s_w_s;
 }
 @end