diff --git a/clang/test/CodeGenCUDA/amdgpu-alias-undef-symbols.cu b/clang/test/CodeGenCUDA/amdgpu-alias-undef-symbols.cu --- a/clang/test/CodeGenCUDA/amdgpu-alias-undef-symbols.cu +++ b/clang/test/CodeGenCUDA/amdgpu-alias-undef-symbols.cu @@ -1,15 +1,14 @@ // 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); }; diff --git a/clang/test/CodeGenCUDA/device-use-host-var.cu b/clang/test/CodeGenCUDA/device-use-host-var.cu --- a/clang/test/CodeGenCUDA/device-use-host-var.cu +++ b/clang/test/CodeGenCUDA/device-use-host-var.cu @@ -1,6 +1,6 @@ -// 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" @@ -71,17 +71,17 @@ // 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; diff --git a/clang/test/CodeGenCUDA/managed-var.cu b/clang/test/CodeGenCUDA/managed-var.cu --- a/clang/test/CodeGenCUDA/managed-var.cu +++ b/clang/test/CodeGenCUDA/managed-var.cu @@ -1,16 +1,16 @@ -// 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 @@ -26,38 +26,38 @@ }; // 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; @@ -81,70 +81,68 @@ } // 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() { @@ -152,18 +150,18 @@ } // 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) diff --git a/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu b/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu --- a/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu +++ b/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu @@ -1,18 +1,18 @@ // 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" @@ -124,9 +124,9 @@ 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 diff --git a/clang/test/CodeGenCUDA/surface.cu b/clang/test/CodeGenCUDA/surface.cu --- a/clang/test/CodeGenCUDA/surface.cu +++ b/clang/test/CodeGenCUDA/surface.cu @@ -1,9 +1,9 @@ // 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; @@ -28,7 +28,7 @@ __attribute__((device)) int suld_2d_zero(surface, 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); @@ -36,7 +36,7 @@ // 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} diff --git a/clang/test/CodeGenCUDA/usual-deallocators.cu b/clang/test/CodeGenCUDA/usual-deallocators.cu --- a/clang/test/CodeGenCUDA/usual-deallocators.cu +++ b/clang/test/CodeGenCUDA/usual-deallocators.cu @@ -1,10 +1,10 @@ -// 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" @@ -83,28 +83,28 @@ // COMMON: call void @_ZN4H1D1dlEPv test_hd(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(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(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(t); // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1D1D2EvPv - // COMMON: call void @_ZN6H1D1D2dlEPv(i8* noundef %3) + // COMMON: call void @_ZN6H1D1D2dlEPv(ptr noundef {{.*}}) test_hd(t); // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1H2D1EvPv - // COMMON: call void @_ZN6H1H2D1dlEPv(i8* {{.*}}) + // COMMON: call void @_ZN6H1H2D1dlEPv(ptr {{.*}}) test_hd(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(t); // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI8H1H2D1D2EvPv - // COMMON: call void @_ZN8H1H2D1D2dlEPv(i8* {{.*}}) + // COMMON: call void @_ZN8H1H2D1D2dlEPv(ptr {{.*}}) test_hd(t); } @@ -113,39 +113,39 @@ // 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} diff --git a/clang/test/CodeGenCUDA/vtbl.cu b/clang/test/CodeGenCUDA/vtbl.cu --- a/clang/test/CodeGenCUDA/vtbl.cu +++ b/clang/test/CodeGenCUDA/vtbl.cu @@ -1,13 +1,12 @@ -// 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() {} }; diff --git a/clang/test/CodeGenCoroutines/coro-params.cpp b/clang/test/CodeGenCoroutines/coro-params.cpp --- a/clang/test/CodeGenCoroutines/coro-params.cpp +++ b/clang/test/CodeGenCoroutines/coro-params.cpp @@ -2,7 +2,7 @@ // 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 struct coroutine_traits; @@ -62,27 +62,25 @@ 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: 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::promise_type"* %__promise to i8* - // CHECK-NEXT: call void @llvm.lifetime.start.p0i8( + // CHECK: store i32 %val, ptr %[[ValAddr:.+]] + + // 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); @@ -93,39 +91,31 @@ // CHECK: call void @_ZN14suspend_always12await_resumeEv( // Destroy promise, then parameter copies: - // CHECK: call void @_ZNSt16coroutine_traitsIJvi8MoveOnly11MoveAndCopyEE12promise_typeD1Ev(%"struct.std::coroutine_traits::promise_type"* {{[^,]*}} %__promise) - // CHECK-NEXT: bitcast %"struct.std::coroutine_traits::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 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::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; @@ -169,10 +159,10 @@ // 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::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; } @@ -195,8 +185,8 @@ 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::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; } diff --git a/clang/test/CodeGenHIP/maybe_undef-attr-verify.hip b/clang/test/CodeGenHIP/maybe_undef-attr-verify.hip --- a/clang/test/CodeGenHIP/maybe_undef-attr-verify.hip +++ b/clang/test/CodeGenHIP/maybe_undef-attr-verify.hip @@ -1,16 +1,16 @@ -// 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:%.*]]) diff --git a/clang/test/CodeGenOpenCL/amdgpu-alignment.cl b/clang/test/CodeGenOpenCL/amdgpu-alignment.cl --- a/clang/test/CodeGenOpenCL/amdgpu-alignment.cl +++ b/clang/test/CodeGenOpenCL/amdgpu-alignment.cl @@ -1,6 +1,6 @@ // 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 @@ -92,48 +92,48 @@ // 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> , <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> , <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> , <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> , <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> , <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> , <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> , <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> , 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> , 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> , 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> , 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> , 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> , 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> , 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]; @@ -379,48 +379,48 @@ // 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> , <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> , <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> , <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> , <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> , <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> , <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> , <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> , 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> , 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> , 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> , 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> , 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> , 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> , 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]; diff --git a/clang/test/CodeGenOpenCL/atomic-ops-libcall.cl b/clang/test/CodeGenOpenCL/atomic-ops-libcall.cl --- a/clang/test/CodeGenOpenCL/atomic-ops-libcall.cl +++ b/clang/test/CodeGenOpenCL/atomic-ops-libcall.cl @@ -1,5 +1,5 @@ -// 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, @@ -20,63 +20,63 @@ 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); } diff --git a/clang/test/CodeGenOpenCL/atomic-ops.cl b/clang/test/CodeGenOpenCL/atomic-ops.cl --- a/clang/test/CodeGenOpenCL/atomic-ops.cl +++ b/clang/test/CodeGenOpenCL/atomic-ops.cl @@ -1,9 +1,9 @@ -// 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 @@ -37,58 +37,58 @@ 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]+]] @@ -105,16 +105,16 @@ // 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); @@ -146,35 +146,35 @@ // 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); } @@ -186,19 +186,19 @@ 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); } @@ -215,10 +215,10 @@ // 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); } @@ -330,13 +330,13 @@ 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); } diff --git a/clang/test/CodeGenOpenCL/blocks.cl b/clang/test/CodeGenOpenCL/blocks.cl --- a/clang/test/CodeGenOpenCL/blocks.cl +++ b/clang/test/CodeGenOpenCL/blocks.cl @@ -1,20 +1,18 @@ -// 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 +// 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: %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*) } +// 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; }; @@ -26,36 +24,32 @@ // 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; @@ -63,14 +57,12 @@ 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 diff --git a/clang/test/CodeGenOpenCL/builtins.cl b/clang/test/CodeGenOpenCL/builtins.cl --- a/clang/test/CodeGenOpenCL/builtins.cl +++ b/clang/test/CodeGenOpenCL/builtins.cl @@ -1,4 +1,4 @@ -// 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. @@ -61,23 +61,20 @@ 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]] } diff --git a/clang/test/CodeGenOpenCL/cast_image.cl b/clang/test/CodeGenOpenCL/cast_image.cl --- a/clang/test/CodeGenOpenCL/cast_image.cl +++ b/clang/test/CodeGenOpenCL/cast_image.cl @@ -1,17 +1,17 @@ -// 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*); } diff --git a/clang/test/CodeGenOpenCL/const-str-array-decay.cl b/clang/test/CodeGenOpenCL/const-str-array-decay.cl --- a/clang/test/CodeGenOpenCL/const-str-array-decay.cl +++ b/clang/test/CodeGenOpenCL/const-str-array-decay.cl @@ -1,4 +1,4 @@ -// 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); @@ -6,6 +6,4 @@ 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 @{{.*}}) diff --git a/clang/test/CodeGenOpenCL/kernels-have-spir-cc-by-default.cl b/clang/test/CodeGenOpenCL/kernels-have-spir-cc-by-default.cl --- a/clang/test/CodeGenOpenCL/kernels-have-spir-cc-by-default.cl +++ b/clang/test/CodeGenOpenCL/kernels-have-spir-cc-by-default.cl @@ -1,5 +1,5 @@ -// 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(). @@ -27,16 +27,16 @@ 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; } @@ -44,8 +44,8 @@ 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; @@ -59,7 +59,7 @@ 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; } diff --git a/clang/test/CodeGenOpenCL/no-half.cl b/clang/test/CodeGenOpenCL/no-half.cl --- a/clang/test/CodeGenOpenCL/no-half.cl +++ b/clang/test/CodeGenOpenCL/no-half.cl @@ -1,39 +1,39 @@ -// 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 } diff --git a/clang/test/CodeGenOpenCL/pipe_builtin.cl b/clang/test/CodeGenOpenCL/pipe_builtin.cl --- a/clang/test/CodeGenOpenCL/pipe_builtin.cl +++ b/clang/test/CodeGenOpenCL/pipe_builtin.cl @@ -1,73 +1,69 @@ -// 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); } diff --git a/clang/test/CodeGenOpenCL/pipe_types.cl b/clang/test/CodeGenOpenCL/pipe_types.cl --- a/clang/test/CodeGenOpenCL/pipe_types.cl +++ b/clang/test/CodeGenOpenCL/pipe_types.cl @@ -1,39 +1,37 @@ -// 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 - -// CHECK: %opencl.pipe_ro_t = type opaque -// CHECK: %opencl.pipe_wo_t = type opaque +// 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 + 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 { @@ -46,7 +44,7 @@ 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) } diff --git a/clang/test/CodeGenOpenCL/to_addr_builtin.cl b/clang/test/CodeGenOpenCL/to_addr_builtin.cl --- a/clang/test/CodeGenOpenCL/to_addr_builtin.cl +++ b/clang/test/CodeGenOpenCL/to_addr_builtin.cl @@ -1,8 +1,7 @@ -// 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; @@ -15,75 +14,75 @@ 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; diff --git a/clang/test/CodeGenSYCL/field-annotate-addr-space.cpp b/clang/test/CodeGenSYCL/field-annotate-addr-space.cpp --- a/clang/test/CodeGenSYCL/field-annotate-addr-space.cpp +++ b/clang/test/CodeGenSYCL/field-annotate-addr-space.cpp @@ -1,4 +1,4 @@ -// 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" @@ -11,9 +11,8 @@ 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; } diff --git a/clang/test/CodeGenSYCL/unique_stable_name.cpp b/clang/test/CodeGenSYCL/unique_stable_name.cpp --- a/clang/test/CodeGenSYCL/unique_stable_name.cpp +++ b/clang/test/CodeGenSYCL/unique_stable_name.cpp @@ -1,4 +1,4 @@ -// 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", @@ -67,48 +67,48 @@ int main() { kernel_single_task(func); - // 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(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( []() { 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(); // CHECK: call spir_func void @_Z14template_paramIiEvv @@ -138,22 +138,22 @@ } // 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))) diff --git a/clang/test/PCH/arc-blocks.mm b/clang/test/PCH/arc-blocks.mm --- a/clang/test/PCH/arc-blocks.mm +++ b/clang/test/PCH/arc-blocks.mm @@ -1,5 +1,5 @@ -// 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 @@ -25,8 +25,6 @@ #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( @@ -36,10 +34,8 @@ } // 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); diff --git a/clang/test/Profile/objc-general.m b/clang/test/Profile/objc-general.m --- a/clang/test/Profile/objc-general.m +++ b/clang/test/Profile/objc-general.m @@ -1,9 +1,9 @@ // 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 @@ -44,7 +44,7 @@ @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; @@ -54,7 +54,7 @@ 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 diff --git a/clang/test/Rewriter/rewrite-super-message.mm b/clang/test/Rewriter/rewrite-super-message.mm --- a/clang/test/Rewriter/rewrite-super-message.mm +++ b/clang/test/Rewriter/rewrite-super-message.mm @@ -1,5 +1,5 @@ -// 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 *); @@ -17,7 +17,7 @@ } @end -// CHECK: call %struct.objc_class* @class_getSuperclass +// CHECK: call ptr @class_getSuperclass @class NSZone; diff --git a/clang/test/SemaObjC/arc-objc-lifetime-conflict.m b/clang/test/SemaObjC/arc-objc-lifetime-conflict.m --- a/clang/test/SemaObjC/arc-objc-lifetime-conflict.m +++ b/clang/test/SemaObjC/arc-objc-lifetime-conflict.m @@ -1,19 +1,12 @@ -// 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