shylie/llvm-project
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[clang] Convert a few tests to opaque pointers

Browse Source 
Reviewed By: nikic

Differential Revision: https://reviews.llvm.org/D150520
This commit is contained in:
Sergei Barannikov 2023-05-14 20:59:13 +03:00
parent 77df976a12
commit f46b0e6d75
27 changed files with 515 additions and 561 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
clang/test/CodeGenCUDA/amdgpu-alias-undef-symbols.cu
View File

@ -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);
};

20
clang/test/CodeGenCUDA/device-use-host-var.cu
View File

@ -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 @@ const int var_host_only = 7;
// 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;

116
clang/test/CodeGenCUDA/managed-var.cu
View File

@ -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 @@ struct vec {
};
// 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 @@ int foo2() {
}
// 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 @@ float addr_taken2() {
}
// 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)

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

@ -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 @@ void foo(const int **a) {
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

10
clang/test/CodeGenCUDA/surface.cu
View File

@ -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 @@ surface<void, 2> surf;
__attribute__((device)) int suld_2d_zero(surface<void, 2>, int, int) asm("llvm.nvvm.suld.2d.i32.zero");
// DEVICE-LABEL: i32 @_Z3fooii(i32 noundef %x, i32 noundef %y)
// DEVICE: call i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)* @surf)
// DEVICE: call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @surf)
// DEVICE: call noundef i32 @llvm.nvvm.suld.2d.i32.zero(i64 %{{.*}}, i32 noundef %{{.*}}, i32 noundef %{{.*}})
__attribute__((device)) int foo(int x, int y) {
return suld_2d_zero(surf, x, y);
@ -36,7 +36,7 @@ __attribute__((device)) int foo(int x, int y) {
// HOST: define internal void @[[PREFIX:__cuda]]_register_globals
// Texture references need registering with correct arguments.
// HOST: call void @[[PREFIX]]RegisterSurface(i8** %0, i8*{{.*}}({{.*}}@surf{{.*}}), i8*{{.*}}({{.*}}@0{{.*}}), i8*{{.*}}({{.*}}@0{{.*}}), i32 2, i32 0)
// HOST: call void @[[PREFIX]]RegisterSurface(ptr %0, ptr @surf, ptr @0, ptr @0, i32 2, i32 0)
// They also need annotating in metadata.
// DEVICE: !0 = !{i64 addrspace(1)* @surf, !"surface", i32 1}
// DEVICE: !0 = !{ptr addrspace(1) @surf, !"surface", i32 1}

52
clang/test/CodeGenCUDA/usual-deallocators.cu
View File

@ -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 @@ __host__ __device__ void tests_hd(void *t) {
// COMMON: call void @_ZN4H1D1dlEPv
test_hd<H1D1>(t);
// COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H1D2EvPv
// DEVICE: call void @_ZN4H1D2dlEPvj(i8* noundef {{.*}}, i32 noundef 1)
// HOST: call void @_ZN4H1D2dlEPv(i8* noundef {{.*}})
// DEVICE: call void @_ZN4H1D2dlEPvj(ptr noundef {{.*}}, i32 noundef 1)
// HOST: call void @_ZN4H1D2dlEPv(ptr noundef {{.*}})
test_hd<H1D2>(t);
// COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H2D1EvPv
// DEVICE: call void @_ZN4H2D1dlEPv(i8* {{.*}})
// HOST: call void @_ZN4H2D1dlEPvj(i8* noundef %3, i32 noundef 1)
// DEVICE: call void @_ZN4H2D1dlEPv(ptr {{.*}})
// HOST: call void @_ZN4H2D1dlEPvj(ptr noundef {{.*}}, i32 noundef 1)
test_hd<H2D1>(t);
// COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H2D2EvPv
// COMMON: call void @_ZN4H2D2dlEPvj(i8* noundef {{.*}}, i32 noundef 1)
// COMMON: call void @_ZN4H2D2dlEPvj(ptr noundef {{.*}}, i32 noundef 1)
test_hd<H2D2>(t);
// COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1D1D2EvPv
// COMMON: call void @_ZN6H1D1D2dlEPv(i8* noundef %3)
// COMMON: call void @_ZN6H1D1D2dlEPv(ptr noundef {{.*}})
test_hd<H1D1D2>(t);
// COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1H2D1EvPv
// COMMON: call void @_ZN6H1H2D1dlEPv(i8* {{.*}})
// COMMON: call void @_ZN6H1H2D1dlEPv(ptr {{.*}})
test_hd<H1H2D1>(t);
// COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1H2D2EvPv
// DEVICE: call void @_ZN6H1H2D2dlEPvj(i8* noundef {{.*}}, i32 noundef 1)
// HOST: call void @_ZN6H1H2D2dlEPv(i8* noundef {{.*}})
// DEVICE: call void @_ZN6H1H2D2dlEPvj(ptr noundef {{.*}}, i32 noundef 1)
// HOST: call void @_ZN6H1H2D2dlEPv(ptr noundef {{.*}})
test_hd<H1H2D2>(t);
// COMMON-LABEL: define linkonce_odr void @_Z7test_hdI8H1H2D1D2EvPv
// COMMON: call void @_ZN8H1H2D1D2dlEPv(i8* {{.*}})
// COMMON: call void @_ZN8H1H2D1D2dlEPv(ptr {{.*}})
test_hd<H1H2D1D2>(t);
}
@ -113,39 +113,39 @@ __host__ __device__ void tests_hd(void *t) {
// 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}

11
clang/test/CodeGenCUDA/vtbl.cu
View File

@ -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() {}
};

82
clang/test/CodeGenCoroutines/coro-params.cpp
View File

@ -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 <typename... T> struct coroutine_traits;
@ -62,27 +62,25 @@ struct MoveAndCopy {
void consume(int,int,int) noexcept;
// TODO: Add support for CopyOnly params
// CHECK: define{{.*}} void @_Z1fi8MoveOnly11MoveAndCopy(i32 noundef %val, %struct.MoveOnly* noundef %[[MoParam:.+]], %struct.MoveAndCopy* noundef %[[McParam:.+]]) #0 personality i8* bitcast (i32 (...)* @__gxx_personality_v0 to i8*
// CHECK: define{{.*}} void @_Z1fi8MoveOnly11MoveAndCopy(i32 noundef %val, ptr noundef %[[MoParam:.+]], ptr noundef %[[McParam:.+]]) #0 personality ptr @__gxx_personality_v0
void f(int val, MoveOnly moParam, MoveAndCopy mcParam) {
// CHECK: %[[MoCopy:.+]] = alloca %struct.MoveOnly,
// CHECK: %[[McCopy:.+]] = alloca %struct.MoveAndCopy,
// CHECK: store i32 %val, i32* %[[ValAddr:.+]]
// CHECK: store i32 %val, ptr %[[ValAddr:.+]]
// CHECK: call i8* @llvm.coro.begin(
// CHECK: call void @_ZN8MoveOnlyC1EOS_(%struct.MoveOnly* {{[^,]*}} %[[MoCopy]], %struct.MoveOnly* noundef nonnull align 4 dereferenceable(4) %[[MoParam]])
// CHECK-NEXT: bitcast %struct.MoveAndCopy* %[[McCopy]] to i8*
// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(
// CHECK-NEXT: call void @_ZN11MoveAndCopyC1EOS_(%struct.MoveAndCopy* {{[^,]*}} %[[McCopy]], %struct.MoveAndCopy* noundef nonnull align 4 dereferenceable(4) %[[McParam]]) #
// CHECK-NEXT: bitcast %"struct.std::coroutine_traits<void, int, MoveOnly, MoveAndCopy>::promise_type"* %__promise to i8*
// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(
// CHECK: call ptr @llvm.coro.begin(
// CHECK: call void @_ZN8MoveOnlyC1EOS_(ptr {{[^,]*}} %[[MoCopy]], ptr noundef nonnull align 4 dereferenceable(4) %[[MoParam]])
// CHECK-NEXT: call void @llvm.lifetime.start.p0(
// CHECK-NEXT: call void @_ZN11MoveAndCopyC1EOS_(ptr {{[^,]*}} %[[McCopy]], ptr noundef nonnull align 4 dereferenceable(4) %[[McParam]]) #
// CHECK-NEXT: call void @llvm.lifetime.start.p0(
// CHECK-NEXT: invoke void @_ZNSt16coroutine_traitsIJvi8MoveOnly11MoveAndCopyEE12promise_typeC1Ev(
// CHECK: call void @_ZN14suspend_always12await_resumeEv(
// CHECK: %[[IntParam:.+]] = load i32, i32* %{{.*}}
// CHECK: %[[MoGep:.+]] = getelementptr inbounds %struct.MoveOnly, %struct.MoveOnly* %[[MoCopy]], i32 0, i32 0
// CHECK: %[[MoVal:.+]] = load i32, i32* %[[MoGep]]
// CHECK: %[[McGep:.+]] = getelementptr inbounds %struct.MoveAndCopy, %struct.MoveAndCopy* %[[McCopy]], i32 0, i32 0
// CHECK: %[[McVal:.+]] = load i32, i32* %[[McGep]]
// CHECK: %[[IntParam:.+]] = load i32, ptr %{{.*}}
// CHECK: %[[MoGep:.+]] = getelementptr inbounds %struct.MoveOnly, ptr %[[MoCopy]], i32 0, i32 0
// CHECK: %[[MoVal:.+]] = load i32, ptr %[[MoGep]]
// CHECK: %[[McGep:.+]] = getelementptr inbounds %struct.MoveAndCopy, ptr %[[McCopy]], i32 0, i32 0
// CHECK: %[[McVal:.+]] = load i32, ptr %[[McGep]]
// CHECK: call void @_Z7consumeiii(i32 noundef %[[IntParam]], i32 noundef %[[MoVal]], i32 noundef %[[McVal]])
consume(val, moParam.val, mcParam.val);
@ -93,39 +91,31 @@ void f(int val, MoveOnly moParam, MoveAndCopy mcParam) {
// CHECK: call void @_ZN14suspend_always12await_resumeEv(
// Destroy promise, then parameter copies:
// CHECK: call void @_ZNSt16coroutine_traitsIJvi8MoveOnly11MoveAndCopyEE12promise_typeD1Ev(%"struct.std::coroutine_traits<void, int, MoveOnly, MoveAndCopy>::promise_type"* {{[^,]*}} %__promise)
// CHECK-NEXT: bitcast %"struct.std::coroutine_traits<void, int, MoveOnly, MoveAndCopy>::promise_type"* %__promise to i8*
// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(
// CHECK-NEXT: call void @_ZN11MoveAndCopyD1Ev(%struct.MoveAndCopy* {{[^,]*}} %[[McCopy]])
// CHECK-NEXT: bitcast %struct.MoveAndCopy* %[[McCopy]] to i8*
// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(
// CHECK-NEXT: call void @_ZN8MoveOnlyD1Ev(%struct.MoveOnly* {{[^,]*}} %[[MoCopy]]
// CHECK-NEXT: bitcast %struct.MoveOnly* %[[MoCopy]] to i8*
// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(
// CHECK-NEXT: bitcast i32* %{{.+}} to i8*
// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(
// CHECK-NEXT: call i8* @llvm.coro.free(
// CHECK: call void @_ZNSt16coroutine_traitsIJvi8MoveOnly11MoveAndCopyEE12promise_typeD1Ev(ptr {{[^,]*}} %__promise)
// CHECK-NEXT: call void @llvm.lifetime.end.p0(
// CHECK-NEXT: call void @_ZN11MoveAndCopyD1Ev(ptr {{[^,]*}} %[[McCopy]])
// CHECK-NEXT: call void @llvm.lifetime.end.p0(
// CHECK-NEXT: call void @_ZN8MoveOnlyD1Ev(ptr {{[^,]*}} %[[MoCopy]]
// CHECK-NEXT: call void @llvm.lifetime.end.p0(
// CHECK-NEXT: call void @llvm.lifetime.end.p0(
// CHECK-NEXT: call ptr @llvm.coro.free(
}
// CHECK-LABEL: void @_Z16dependent_paramsI1A1BEvT_T0_S3_(%struct.A* noundef %x, %struct.B* noundef %0, %struct.B* noundef %y)
// CHECK-LABEL: void @_Z16dependent_paramsI1A1BEvT_T0_S3_(ptr noundef %x, ptr noundef %0, ptr noundef %y)
template <typename T, typename U>
void dependent_params(T x, U, U y) {
// CHECK: %[[x_copy:.+]] = alloca %struct.A,
// CHECK-NEXT: %[[unnamed_copy:.+]] = alloca %struct.B
// CHECK-NEXT: %[[y_copy:.+]] = alloca %struct.B
// CHECK: call i8* @llvm.coro.begin
// CHECK-NEXT: bitcast %struct.A* %[[x_copy]] to i8*
// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(
// CHECK-NEXT: call void @_ZN1AC1EOS_(%struct.A* {{[^,]*}} %[[x_copy]], %struct.A* noundef nonnull align 4 dereferenceable(512) %x)
// CHECK-NEXT: bitcast %struct.B* %[[unnamed_copy]] to i8*
// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(
// CHECK-NEXT: call void @_ZN1BC1EOS_(%struct.B* {{[^,]*}} %[[unnamed_copy]], %struct.B* noundef nonnull align 4 dereferenceable(512) %0)
// CHECK-NEXT: bitcast %struct.B* %[[y_copy]] to i8*
// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(
// CHECK-NEXT: call void @_ZN1BC1EOS_(%struct.B* {{[^,]*}} %[[y_copy]], %struct.B* noundef nonnull align 4 dereferenceable(512) %y)
// CHECK-NEXT: bitcast %"struct.std::coroutine_traits<void, A, B, B>::promise_type"* %__promise to i8*
// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(
// CHECK: call ptr @llvm.coro.begin
// CHECK-NEXT: call void @llvm.lifetime.start.p0(
// CHECK-NEXT: call void @_ZN1AC1EOS_(ptr {{[^,]*}} %[[x_copy]], ptr noundef nonnull align 4 dereferenceable(512) %x)
// CHECK-NEXT: call void @llvm.lifetime.start.p0(
// CHECK-NEXT: call void @_ZN1BC1EOS_(ptr {{[^,]*}} %[[unnamed_copy]], ptr noundef nonnull align 4 dereferenceable(512) %0)
// CHECK-NEXT: call void @llvm.lifetime.start.p0(
// CHECK-NEXT: call void @_ZN1BC1EOS_(ptr {{[^,]*}} %[[y_copy]], ptr noundef nonnull align 4 dereferenceable(512) %y)
// CHECK-NEXT: call void @llvm.lifetime.start.p0(
// CHECK-NEXT: invoke void @_ZNSt16coroutine_traitsIJv1A1BS1_EE12promise_typeC1Ev(
co_return;
@ -169,10 +159,10 @@ struct std::coroutine_traits<void, promise_matching_constructor, int, float, dou
// CHECK-LABEL: void @_Z38coroutine_matching_promise_constructor28promise_matching_constructorifd(i32 noundef %0, float noundef %1, double noundef %2)
void coroutine_matching_promise_constructor(promise_matching_constructor, int, float, double) {
// CHECK: %[[INT:.+]] = load i32, i32* %5, align 4
// CHECK: %[[FLOAT:.+]] = load float, float* %6, align 4
// CHECK: %[[DOUBLE:.+]] = load double, double* %7, align 8
// CHECK: invoke void @_ZNSt16coroutine_traitsIJv28promise_matching_constructorifdEE12promise_typeC1ES0_ifd(%"struct.std::coroutine_traits<void, promise_matching_constructor, int, float, double>::promise_type"* {{[^,]*}} %__promise, i32 noundef %[[INT]], float noundef %[[FLOAT]], double noundef %[[DOUBLE]])
// CHECK: %[[INT:.+]] = load i32, ptr %5, align 4
// CHECK: %[[FLOAT:.+]] = load float, ptr %6, align 4
// CHECK: %[[DOUBLE:.+]] = load double, ptr %7, align 8
// CHECK: invoke void @_ZNSt16coroutine_traitsIJv28promise_matching_constructorifdEE12promise_typeC1ES0_ifd(ptr {{[^,]*}} %__promise, i32 noundef %[[INT]], float noundef %[[FLOAT]], double noundef %[[DOUBLE]])
co_return;
}
@ -195,8 +185,8 @@ struct some_class {
method good_coroutine_calls_custom_constructor(float);
};
// CHECK-LABEL: define{{.*}} void @_ZN10some_class39good_coroutine_calls_custom_constructorEf(%struct.some_class*
// CHECK-LABEL: define{{.*}} void @_ZN10some_class39good_coroutine_calls_custom_constructorEf(ptr
method some_class::good_coroutine_calls_custom_constructor(float) {
// CHECK: invoke void @_ZNSt16coroutine_traitsIJ6methodR10some_classfEE12promise_typeC1ES2_f(%"struct.std::coroutine_traits<method, some_class &, float>::promise_type"* {{[^,]*}} %__promise, %struct.some_class* noundef nonnull align 1 dereferenceable(1) %{{.+}}, float
// CHECK: invoke void @_ZNSt16coroutine_traitsIJ6methodR10some_classfEE12promise_typeC1ES2_f(ptr {{[^,]*}} %__promise, ptr noundef nonnull align 1 dereferenceable(1) %{{.+}}, float
co_return;
}

10
clang/test/CodeGenHIP/maybe_undef-attr-verify.hip
View File

@ -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:%.*]])

172
clang/test/CodeGenOpenCL/amdgpu-alignment.cl
View File

@ -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 @@ typedef double __attribute__((ext_vector_type(16))) double16;
// CHECK-LABEL: @local_memory_alignment_global(
// CHECK: store volatile i8 0, i8 addrspace(3)* getelementptr inbounds ([4 x i8], [4 x i8] addrspace(3)* @local_memory_alignment_global.lds_i8, i64 0, i64 0), align 1
// CHECK: store volatile <2 x i8> zeroinitializer, <2 x i8> addrspace(3)* getelementptr inbounds ([4 x <2 x i8>], [4 x <2 x i8>] addrspace(3)* @local_memory_alignment_global.lds_v2i8, i64 0, i64 0), align 2
// CHECK: store volatile <4 x i8> <i8 0, i8 0, i8 0, i8 undef>, <4 x i8> addrspace(3)* bitcast ([4 x <3 x i8>] addrspace(3)* @local_memory_alignment_global.lds_v3i8 to <4 x i8> addrspace(3)*), align 4
// CHECK: store volatile <4 x i8> zeroinitializer, <4 x i8> addrspace(3)* getelementptr inbounds ([4 x <4 x i8>], [4 x <4 x i8>] addrspace(3)* @local_memory_alignment_global.lds_v4i8, i64 0, i64 0), align 4
// CHECK: store volatile <8 x i8> zeroinitializer, <8 x i8> addrspace(3)* getelementptr inbounds ([4 x <8 x i8>], [4 x <8 x i8>] addrspace(3)* @local_memory_alignment_global.lds_v8i8, i64 0, i64 0), align 8
// CHECK: store volatile <16 x i8> zeroinitializer, <16 x i8> addrspace(3)* getelementptr inbounds ([4 x <16 x i8>], [4 x <16 x i8>] addrspace(3)* @local_memory_alignment_global.lds_v16i8, i64 0, i64 0), align 16
// CHECK: store volatile i16 0, i16 addrspace(3)* getelementptr inbounds ([4 x i16], [4 x i16] addrspace(3)* @local_memory_alignment_global.lds_i16, i64 0, i64 0), align 2
// CHECK: store volatile <2 x i16> zeroinitializer, <2 x i16> addrspace(3)* getelementptr inbounds ([4 x <2 x i16>], [4 x <2 x i16>] addrspace(3)* @local_memory_alignment_global.lds_v2i16, i64 0, i64 0), align 4
// CHECK: store volatile <4 x i16> <i16 0, i16 0, i16 0, i16 undef>, <4 x i16> addrspace(3)* bitcast ([4 x <3 x i16>] addrspace(3)* @local_memory_alignment_global.lds_v3i16 to <4 x i16> addrspace(3)*), align 8
// CHECK: store volatile <4 x i16> zeroinitializer, <4 x i16> addrspace(3)* getelementptr inbounds ([4 x <4 x i16>], [4 x <4 x i16>] addrspace(3)* @local_memory_alignment_global.lds_v4i16, i64 0, i64 0), align 8
// CHECK: store volatile <8 x i16> zeroinitializer, <8 x i16> addrspace(3)* getelementptr inbounds ([4 x <8 x i16>], [4 x <8 x i16>] addrspace(3)* @local_memory_alignment_global.lds_v8i16, i64 0, i64 0), align 16
// CHECK: store volatile <16 x i16> zeroinitializer, <16 x i16> addrspace(3)* getelementptr inbounds ([4 x <16 x i16>], [4 x <16 x i16>] addrspace(3)* @local_memory_alignment_global.lds_v16i16, i64 0, i64 0), align 32
// CHECK: store volatile i32 0, i32 addrspace(3)* getelementptr inbounds ([4 x i32], [4 x i32] addrspace(3)* @local_memory_alignment_global.lds_i32, i64 0, i64 0), align 4
// CHECK: store volatile <2 x i32> zeroinitializer, <2 x i32> addrspace(3)* getelementptr inbounds ([4 x <2 x i32>], [4 x <2 x i32>] addrspace(3)* @local_memory_alignment_global.lds_v2i32, i64 0, i64 0), align 8
// CHECK: store volatile <4 x i32> <i32 0, i32 0, i32 0, i32 undef>, <4 x i32> addrspace(3)* bitcast ([4 x <3 x i32>] addrspace(3)* @local_memory_alignment_global.lds_v3i32 to <4 x i32> addrspace(3)*), align 16
// CHECK: store volatile <4 x i32> zeroinitializer, <4 x i32> addrspace(3)* getelementptr inbounds ([4 x <4 x i32>], [4 x <4 x i32>] addrspace(3)* @local_memory_alignment_global.lds_v4i32, i64 0, i64 0), align 16
// CHECK: store volatile <8 x i32> zeroinitializer, <8 x i32> addrspace(3)* getelementptr inbounds ([4 x <8 x i32>], [4 x <8 x i32>] addrspace(3)* @local_memory_alignment_global.lds_v8i32, i64 0, i64 0), align 32
// CHECK: store volatile <16 x i32> zeroinitializer, <16 x i32> addrspace(3)* getelementptr inbounds ([4 x <16 x i32>], [4 x <16 x i32>] addrspace(3)* @local_memory_alignment_global.lds_v16i32, i64 0, i64 0), align 64
// CHECK: store volatile i64 0, i64 addrspace(3)* getelementptr inbounds ([4 x i64], [4 x i64] addrspace(3)* @local_memory_alignment_global.lds_i64, i64 0, i64 0), align 8
// CHECK: store volatile <2 x i64> zeroinitializer, <2 x i64> addrspace(3)* getelementptr inbounds ([4 x <2 x i64>], [4 x <2 x i64>] addrspace(3)* @local_memory_alignment_global.lds_v2i64, i64 0, i64 0), align 16
// CHECK: store volatile <4 x i64> <i64 0, i64 0, i64 0, i64 undef>, <4 x i64> addrspace(3)* bitcast ([4 x <3 x i64>] addrspace(3)* @local_memory_alignment_global.lds_v3i64 to <4 x i64> addrspace(3)*), align 32
// CHECK: store volatile <4 x i64> zeroinitializer, <4 x i64> addrspace(3)* getelementptr inbounds ([4 x <4 x i64>], [4 x <4 x i64>] addrspace(3)* @local_memory_alignment_global.lds_v4i64, i64 0, i64 0), align 32
// CHECK: store volatile <8 x i64> zeroinitializer, <8 x i64> addrspace(3)* getelementptr inbounds ([4 x <8 x i64>], [4 x <8 x i64>] addrspace(3)* @local_memory_alignment_global.lds_v8i64, i64 0, i64 0), align 64
// CHECK: store volatile <16 x i64> zeroinitializer, <16 x i64> addrspace(3)* getelementptr inbounds ([4 x <16 x i64>], [4 x <16 x i64>] addrspace(3)* @local_memory_alignment_global.lds_v16i64, i64 0, i64 0), align 128
// CHECK: store volatile half 0xH0000, half addrspace(3)* getelementptr inbounds ([4 x half], [4 x half] addrspace(3)* @local_memory_alignment_global.lds_f16, i64 0, i64 0), align 2
// CHECK: store volatile <2 x half> zeroinitializer, <2 x half> addrspace(3)* getelementptr inbounds ([4 x <2 x half>], [4 x <2 x half>] addrspace(3)* @local_memory_alignment_global.lds_v2f16, i64 0, i64 0), align 4
// CHECK: store volatile <4 x half> <half 0xH0000, half 0xH0000, half 0xH0000, half undef>, <4 x half> addrspace(3)* bitcast ([4 x <3 x half>] addrspace(3)* @local_memory_alignment_global.lds_v3f16 to <4 x half> addrspace(3)*), align 8
// CHECK: store volatile <4 x half> zeroinitializer, <4 x half> addrspace(3)* getelementptr inbounds ([4 x <4 x half>], [4 x <4 x half>] addrspace(3)* @local_memory_alignment_global.lds_v4f16, i64 0, i64 0), align 8
// CHECK: store volatile <8 x half> zeroinitializer, <8 x half> addrspace(3)* getelementptr inbounds ([4 x <8 x half>], [4 x <8 x half>] addrspace(3)* @local_memory_alignment_global.lds_v8f16, i64 0, i64 0), align 16
// CHECK: store volatile <16 x half> zeroinitializer, <16 x half> addrspace(3)* getelementptr inbounds ([4 x <16 x half>], [4 x <16 x half>] addrspace(3)* @local_memory_alignment_global.lds_v16f16, i64 0, i64 0), align 32
// CHECK: store volatile float 0.000000e+00, float addrspace(3)* getelementptr inbounds ([4 x float], [4 x float] addrspace(3)* @local_memory_alignment_global.lds_f32, i64 0, i64 0), align 4
// CHECK: store volatile <2 x float> zeroinitializer, <2 x float> addrspace(3)* getelementptr inbounds ([4 x <2 x float>], [4 x <2 x float>] addrspace(3)* @local_memory_alignment_global.lds_v2f32, i64 0, i64 0), align 8
// CHECK: store volatile <4 x float> <float 0.000000e+00, float 0.000000e+00, float 0.000000e+00, float undef>, <4 x float> addrspace(3)* bitcast ([4 x <3 x float>] addrspace(3)* @local_memory_alignment_global.lds_v3f32 to <4 x float> addrspace(3)*), align 16
// CHECK: store volatile <4 x float> zeroinitializer, <4 x float> addrspace(3)* getelementptr inbounds ([4 x <4 x float>], [4 x <4 x float>] addrspace(3)* @local_memory_alignment_global.lds_v4f32, i64 0, i64 0), align 16
// CHECK: store volatile <8 x float> zeroinitializer, <8 x float> addrspace(3)* getelementptr inbounds ([4 x <8 x float>], [4 x <8 x float>] addrspace(3)* @local_memory_alignment_global.lds_v8f32, i64 0, i64 0), align 32
// CHECK: store volatile <16 x float> zeroinitializer, <16 x float> addrspace(3)* getelementptr inbounds ([4 x <16 x float>], [4 x <16 x float>] addrspace(3)* @local_memory_alignment_global.lds_v16f32, i64 0, i64 0), align 64
// CHECK: store volatile double 0.000000e+00, double addrspace(3)* getelementptr inbounds ([4 x double], [4 x double] addrspace(3)* @local_memory_alignment_global.lds_f64, i64 0, i64 0), align 8
// CHECK: store volatile <2 x double> zeroinitializer, <2 x double> addrspace(3)* getelementptr inbounds ([4 x <2 x double>], [4 x <2 x double>] addrspace(3)* @local_memory_alignment_global.lds_v2f64, i64 0, i64 0), align 16
// CHECK: store volatile <4 x double> <double 0.000000e+00, double 0.000000e+00, double 0.000000e+00, double undef>, <4 x double> addrspace(3)* bitcast ([4 x <3 x double>] addrspace(3)* @local_memory_alignment_global.lds_v3f64 to <4 x double> addrspace(3)*), align 32
// CHECK: store volatile <4 x double> zeroinitializer, <4 x double> addrspace(3)* getelementptr inbounds ([4 x <4 x double>], [4 x <4 x double>] addrspace(3)* @local_memory_alignment_global.lds_v4f64, i64 0, i64 0), align 32
// CHECK: store volatile <8 x double> zeroinitializer, <8 x double> addrspace(3)* getelementptr inbounds ([4 x <8 x double>], [4 x <8 x double>] addrspace(3)* @local_memory_alignment_global.lds_v8f64, i64 0, i64 0), align 64
// CHECK: store volatile <16 x double> zeroinitializer, <16 x double> addrspace(3)* getelementptr inbounds ([4 x <16 x double>], [4 x <16 x double>] addrspace(3)* @local_memory_alignment_global.lds_v16f64, i64 0, i64 0), align 128
// CHECK: store volatile i8 0, ptr addrspace(3) @local_memory_alignment_global.lds_i8, align 1
// CHECK: store volatile <2 x i8> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v2i8, align 2
// CHECK: store volatile <4 x i8> <i8 0, i8 0, i8 0, i8 undef>, ptr addrspace(3) @local_memory_alignment_global.lds_v3i8, align 4
// CHECK: store volatile <4 x i8> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v4i8, align 4
// CHECK: store volatile <8 x i8> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v8i8, align 8
// CHECK: store volatile <16 x i8> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v16i8, align 16
// CHECK: store volatile i16 0, ptr addrspace(3) @local_memory_alignment_global.lds_i16, align 2
// CHECK: store volatile <2 x i16> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v2i16, align 4
// CHECK: store volatile <4 x i16> <i16 0, i16 0, i16 0, i16 undef>, ptr addrspace(3) @local_memory_alignment_global.lds_v3i16, align 8
// CHECK: store volatile <4 x i16> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v4i16, align 8
// CHECK: store volatile <8 x i16> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v8i16, align 16
// CHECK: store volatile <16 x i16> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v16i16, align 32
// CHECK: store volatile i32 0, ptr addrspace(3) @local_memory_alignment_global.lds_i32, align 4
// CHECK: store volatile <2 x i32> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v2i32, align 8
// CHECK: store volatile <4 x i32> <i32 0, i32 0, i32 0, i32 undef>, ptr addrspace(3) @local_memory_alignment_global.lds_v3i32, align 16
// CHECK: store volatile <4 x i32> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v4i32, align 16
// CHECK: store volatile <8 x i32> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v8i32, align 32
// CHECK: store volatile <16 x i32> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v16i32, align 64
// CHECK: store volatile i64 0, ptr addrspace(3) @local_memory_alignment_global.lds_i64, align 8
// CHECK: store volatile <2 x i64> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v2i64, align 16
// CHECK: store volatile <4 x i64> <i64 0, i64 0, i64 0, i64 undef>, ptr addrspace(3) @local_memory_alignment_global.lds_v3i64, align 32
// CHECK: store volatile <4 x i64> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v4i64, align 32
// CHECK: store volatile <8 x i64> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v8i64, align 64
// CHECK: store volatile <16 x i64> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v16i64, align 128
// CHECK: store volatile half 0xH0000, ptr addrspace(3) @local_memory_alignment_global.lds_f16, align 2
// CHECK: store volatile <2 x half> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v2f16, align 4
// CHECK: store volatile <4 x half> <half 0xH0000, half 0xH0000, half 0xH0000, half undef>, ptr addrspace(3) @local_memory_alignment_global.lds_v3f16, align 8
// CHECK: store volatile <4 x half> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v4f16, align 8
// CHECK: store volatile <8 x half> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v8f16, align 16
// CHECK: store volatile <16 x half> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v16f16, align 32
// CHECK: store volatile float 0.000000e+00, ptr addrspace(3) @local_memory_alignment_global.lds_f32, align 4
// CHECK: store volatile <2 x float> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v2f32, align 8
// CHECK: store volatile <4 x float> <float 0.000000e+00, float 0.000000e+00, float 0.000000e+00, float undef>, ptr addrspace(3) @local_memory_alignment_global.lds_v3f32, align 16
// CHECK: store volatile <4 x float> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v4f32, align 16
// CHECK: store volatile <8 x float> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v8f32, align 32
// CHECK: store volatile <16 x float> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v16f32, align 64
// CHECK: store volatile double 0.000000e+00, ptr addrspace(3) @local_memory_alignment_global.lds_f64, align 8
// CHECK: store volatile <2 x double> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v2f64, align 16
// CHECK: store volatile <4 x double> <double 0.000000e+00, double 0.000000e+00, double 0.000000e+00, double undef>, ptr addrspace(3) @local_memory_alignment_global.lds_v3f64, align 32
// CHECK: store volatile <4 x double> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v4f64, align 32
// CHECK: store volatile <8 x double> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v8f64, align 64
// CHECK: store volatile <16 x double> zeroinitializer, ptr addrspace(3) @local_memory_alignment_global.lds_v16f64, align 128
kernel void local_memory_alignment_global()
{
volatile local char lds_i8[4];
@ -379,48 +379,48 @@ kernel void local_memory_alignment_arg(
// CHECK: %private_v8f64 = alloca [4 x <8 x double>], align 64, addrspace(5)
// CHECK: %private_v16f64 = alloca [4 x <16 x double>], align 128, addrspace(5)
// CHECK: store volatile i8 0, i8 addrspace(5)* %arraydecay, align 1
// CHECK: store volatile <2 x i8> zeroinitializer, <2 x i8> addrspace(5)* %arraydecay{{[0-9]+}}, align 2
// CHECK: store volatile <4 x i8> <i8 0, i8 0, i8 0, i8 undef>, <4 x i8> addrspace(5)* %storetmp, align 4
// CHECK: store volatile <4 x i8> zeroinitializer, <4 x i8> addrspace(5)* %arraydecay{{[0-9]+}}, align 4
// CHECK: store volatile <8 x i8> zeroinitializer, <8 x i8> addrspace(5)* %arraydecay{{[0-9]+}}, align 8
// CHECK: store volatile <16 x i8> zeroinitializer, <16 x i8> addrspace(5)* %arraydecay{{[0-9]+}}, align 16
// CHECK: store volatile i16 0, i16 addrspace(5)* %arraydecay{{[0-9]+}}, align 2
// CHECK: store volatile <2 x i16> zeroinitializer, <2 x i16> addrspace(5)* %arraydecay{{[0-9]+}}, align 4
// CHECK: store volatile <4 x i16> <i16 0, i16 0, i16 0, i16 undef>, <4 x i16> addrspace(5)* %storetmp{{[0-9]+}}, align 8
// CHECK: store volatile <4 x i16> zeroinitializer, <4 x i16> addrspace(5)* %arraydecay{{[0-9]+}}, align 8
// CHECK: store volatile <8 x i16> zeroinitializer, <8 x i16> addrspace(5)* %arraydecay{{[0-9]+}}, align 16
// CHECK: store volatile <16 x i16> zeroinitializer, <16 x i16> addrspace(5)* %arraydecay{{[0-9]+}}, align 32
// CHECK: store volatile i32 0, i32 addrspace(5)* %arraydecay{{[0-9]+}}, align 4
// CHECK: store volatile <2 x i32> zeroinitializer, <2 x i32> addrspace(5)* %arraydecay{{[0-9]+}}, align 8
// CHECK: store volatile <4 x i32> <i32 0, i32 0, i32 0, i32 undef>, <4 x i32> addrspace(5)* %storetmp16, align 16
// CHECK: store volatile <4 x i32> zeroinitializer, <4 x i32> addrspace(5)* %arraydecay{{[0-9]+}}, align 16
// CHECK: store volatile <8 x i32> zeroinitializer, <8 x i32> addrspace(5)* %arraydecay{{[0-9]+}}, align 32
// CHECK: store volatile <16 x i32> zeroinitializer, <16 x i32> addrspace(5)* %arraydecay{{[0-9]+}}, align 64
// CHECK: store volatile i64 0, i64 addrspace(5)* %arraydecay{{[0-9]+}}, align 8
// CHECK: store volatile <2 x i64> zeroinitializer, <2 x i64> addrspace(5)* %arraydecay{{[0-9]+}}, align 16
// CHECK: store volatile <4 x i64> <i64 0, i64 0, i64 0, i64 undef>, <4 x i64> addrspace(5)* %storetmp23, align 32
// CHECK: store volatile <4 x i64> zeroinitializer, <4 x i64> addrspace(5)* %arraydecay{{[0-9]+}}, align 32
// CHECK: store volatile <8 x i64> zeroinitializer, <8 x i64> addrspace(5)* %arraydecay{{[0-9]+}}, align 64
// CHECK: store volatile <16 x i64> zeroinitializer, <16 x i64> addrspace(5)* %arraydecay{{[0-9]+}}, align 128
// CHECK: store volatile half 0xH0000, half addrspace(5)* %arraydecay{{[0-9]+}}, align 2
// CHECK: store volatile <2 x half> zeroinitializer, <2 x half> addrspace(5)* %arraydecay{{[0-9]+}}, align 4
// CHECK: store volatile <4 x half> <half 0xH0000, half 0xH0000, half 0xH0000, half undef>, <4 x half> addrspace(5)* %storetmp{{[0-9]+}}, align 8
// CHECK: store volatile <4 x half> zeroinitializer, <4 x half> addrspace(5)* %arraydecay{{[0-9]+}}, align 8
// CHECK: store volatile <8 x half> zeroinitializer, <8 x half> addrspace(5)* %arraydecay{{[0-9]+}}, align 16
// CHECK: store volatile <16 x half> zeroinitializer, <16 x half> addrspace(5)* %arraydecay{{[0-9]+}}, align 32
// CHECK: store volatile float 0.000000e+00, float addrspace(5)* %arraydecay34, align 4
// CHECK: store volatile <2 x float> zeroinitializer, <2 x float> addrspace(5)* %arraydecay{{[0-9]+}}, align 8
// CHECK: store volatile <4 x float> <float 0.000000e+00, float 0.000000e+00, float 0.000000e+00, float undef>, <4 x float> addrspace(5)* %storetmp{{[0-9]+}}, align 16
// CHECK: store volatile <4 x float> zeroinitializer, <4 x float> addrspace(5)* %arraydecay{{[0-9]+}}, align 16
// CHECK: store volatile <8 x float> zeroinitializer, <8 x float> addrspace(5)* %arraydecay{{[0-9]+}}, align 32
// CHECK: store volatile <16 x float> zeroinitializer, <16 x float> addrspace(5)* %arraydecay{{[0-9]+}}, align 64
// CHECK: store volatile double 0.000000e+00, double addrspace(5)* %arraydecay{{[0-9]+}}, align 8
// CHECK: store volatile <2 x double> zeroinitializer, <2 x double> addrspace(5)* %arraydecay{{[0-9]+}}, align 16
// CHECK: store volatile <4 x double> <double 0.000000e+00, double 0.000000e+00, double 0.000000e+00, double undef>, <4 x double> addrspace(5)* %storetmp{{[0-9]+}}, align 32
// CHECK: store volatile <4 x double> zeroinitializer, <4 x double> addrspace(5)* %arraydecay{{[0-9]+}}, align 32
// CHECK: store volatile <8 x double> zeroinitializer, <8 x double> addrspace(5)* %arraydecay{{[0-9]+}}, align 64
// CHECK: store volatile <16 x double> zeroinitializer, <16 x double> addrspace(5)* %arraydecay{{[0-9]+}}, align 128
// CHECK: store volatile i8 0, ptr addrspace(5) %arraydecay, align 1
// CHECK: store volatile <2 x i8> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 2
// CHECK: store volatile <4 x i8> <i8 0, i8 0, i8 0, i8 undef>, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 4
// CHECK: store volatile <4 x i8> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 4
// CHECK: store volatile <8 x i8> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 8
// CHECK: store volatile <16 x i8> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 16
// CHECK: store volatile i16 0, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 2
// CHECK: store volatile <2 x i16> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 4
// CHECK: store volatile <4 x i16> <i16 0, i16 0, i16 0, i16 undef>, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 8
// CHECK: store volatile <4 x i16> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 8
// CHECK: store volatile <8 x i16> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 16
// CHECK: store volatile <16 x i16> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 32
// CHECK: store volatile i32 0, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 4
// CHECK: store volatile <2 x i32> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 8
// CHECK: store volatile <4 x i32> <i32 0, i32 0, i32 0, i32 undef>, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 16
// CHECK: store volatile <4 x i32> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 16
// CHECK: store volatile <8 x i32> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 32
// CHECK: store volatile <16 x i32> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 64
// CHECK: store volatile i64 0, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 8
// CHECK: store volatile <2 x i64> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 16
// CHECK: store volatile <4 x i64> <i64 0, i64 0, i64 0, i64 undef>, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 32
// CHECK: store volatile <4 x i64> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 32
// CHECK: store volatile <8 x i64> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 64
// CHECK: store volatile <16 x i64> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 128
// CHECK: store volatile half 0xH0000, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 2
// CHECK: store volatile <2 x half> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 4
// CHECK: store volatile <4 x half> <half 0xH0000, half 0xH0000, half 0xH0000, half undef>, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 8
// CHECK: store volatile <4 x half> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 8
// CHECK: store volatile <8 x half> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 16
// CHECK: store volatile <16 x half> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 32
// CHECK: store volatile float 0.000000e+00, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 4
// CHECK: store volatile <2 x float> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 8
// CHECK: store volatile <4 x float> <float 0.000000e+00, float 0.000000e+00, float 0.000000e+00, float undef>, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 16
// CHECK: store volatile <4 x float> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 16
// CHECK: store volatile <8 x float> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 32
// CHECK: store volatile <16 x float> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 64
// CHECK: store volatile double 0.000000e+00, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 8
// CHECK: store volatile <2 x double> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 16
// CHECK: store volatile <4 x double> <double 0.000000e+00, double 0.000000e+00, double 0.000000e+00, double undef>, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 32
// CHECK: store volatile <4 x double> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 32
// CHECK: store volatile <8 x double> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 64
// CHECK: store volatile <16 x double> zeroinitializer, ptr addrspace(5) %arraydecay{{[0-9]+}}, align 128
kernel void private_memory_alignment_alloca()
{
volatile private char private_i8[4];

64
clang/test/CodeGenOpenCL/atomic-ops-libcall.cl
View File

@ -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 @@ typedef enum memory_scope {
void f(atomic_int *i, global atomic_int *gi, local atomic_int *li, private atomic_int *pi, atomic_uint *ui, int cmp, int order, int scope) {
int x;
// SPIR: {{%[^ ]*}} = call i32 @__opencl_atomic_load_4(i8 addrspace(4)* noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// ARM: {{%[^ ]*}} = call i32 @__opencl_atomic_load_4(i8* noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// SPIR: {{%[^ ]*}} = call i32 @__opencl_atomic_load_4(ptr addrspace(4) noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// ARM: {{%[^ ]*}} = call i32 @__opencl_atomic_load_4(ptr noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
x = __opencl_atomic_load(i, memory_order_seq_cst, memory_scope_work_group);
// SPIR: call void @__opencl_atomic_store_4(i8 addrspace(4)* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// ARM: call void @__opencl_atomic_store_4(i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// SPIR: call void @__opencl_atomic_store_4(ptr addrspace(4) noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// ARM: call void @__opencl_atomic_store_4(ptr noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
__opencl_atomic_store(i, 1, memory_order_seq_cst, memory_scope_work_group);
// SPIR: %[[GP:[0-9]+]] = addrspacecast i8 addrspace(1)* {{%[0-9]+}} to i8 addrspace(4)*
// SPIR: call void @__opencl_atomic_store_4(i8 addrspace(4)* noundef %[[GP]], i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// ARM: call void @__opencl_atomic_store_4(i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// SPIR: %[[GP:[0-9]+]] = addrspacecast ptr addrspace(1) {{%[0-9]+}} to ptr addrspace(4)
// SPIR: call void @__opencl_atomic_store_4(ptr addrspace(4) noundef %[[GP]], i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// ARM: call void @__opencl_atomic_store_4(ptr noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
__opencl_atomic_store(gi, 1, memory_order_seq_cst, memory_scope_work_group);
// SPIR: %[[GP:[0-9]+]] = addrspacecast i8 addrspace(3)* {{%[0-9]+}} to i8 addrspace(4)*
// SPIR: call void @__opencl_atomic_store_4(i8 addrspace(4)* noundef %[[GP]], i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// ARM: call void @__opencl_atomic_store_4(i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// SPIR: %[[GP:[0-9]+]] = addrspacecast ptr addrspace(3) {{%[0-9]+}} to ptr addrspace(4)
// SPIR: call void @__opencl_atomic_store_4(ptr addrspace(4) noundef %[[GP]], i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// ARM: call void @__opencl_atomic_store_4(ptr noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
__opencl_atomic_store(li, 1, memory_order_seq_cst, memory_scope_work_group);
// SPIR: %[[GP:[0-9]+]] = addrspacecast i8* {{%[0-9]+}} to i8 addrspace(4)*
// SPIR: call void @__opencl_atomic_store_4(i8 addrspace(4)* noundef %[[GP]], i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// ARM: call void @__opencl_atomic_store_4(i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// SPIR: %[[GP:[0-9]+]] = addrspacecast ptr {{%[0-9]+}} to ptr addrspace(4)
// SPIR: call void @__opencl_atomic_store_4(ptr addrspace(4) noundef %[[GP]], i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// ARM: call void @__opencl_atomic_store_4(ptr noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
__opencl_atomic_store(pi, 1, memory_order_seq_cst, memory_scope_work_group);
// SPIR: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_add_4(i8 addrspace(4)* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// ARM: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_add_4(i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// SPIR: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_add_4(ptr addrspace(4) noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// ARM: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_add_4(ptr noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
x = __opencl_atomic_fetch_add(i, 3, memory_order_seq_cst, memory_scope_work_group);
// SPIR: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_min_4(i8 addrspace(4)* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// ARM: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_min_4(i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// SPIR: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_min_4(ptr addrspace(4) noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// ARM: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_min_4(ptr noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
x = __opencl_atomic_fetch_min(i, 3, memory_order_seq_cst, memory_scope_work_group);
// SPIR: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_umin_4(i8 addrspace(4)* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// ARM: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_umin_4(i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// SPIR: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_umin_4(ptr addrspace(4) noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
// ARM: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_umin_4(ptr noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 1)
x = __opencl_atomic_fetch_min(ui, 3, memory_order_seq_cst, memory_scope_work_group);
// SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* noundef {{%[0-9]+}}, i8 addrspace(4)* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 1)
// ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8* noundef {{%[0-9]+}}, i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 1)
// SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(ptr addrspace(4) noundef {{%[0-9]+}}, ptr addrspace(4) noundef {{%[^,]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 1)
// ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(ptr noundef {{%[0-9]+}}, ptr noundef {{%[^,]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 1)
x = __opencl_atomic_compare_exchange_strong(i, &cmp, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
// SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* noundef {{%[0-9]+}}, i8 addrspace(4)* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 1)
// ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8* noundef {{%[0-9]+}}, i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 1)
// SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(ptr addrspace(4) noundef {{%[0-9]+}}, ptr addrspace(4) noundef {{%[^,]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 1)
// ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(ptr noundef {{%[0-9]+}}, ptr noundef {{%[^,]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 1)
x = __opencl_atomic_compare_exchange_weak(i, &cmp, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
// SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* noundef {{%[0-9]+}}, i8 addrspace(4)* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 2)
// ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8* noundef {{%[0-9]+}}, i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 2)
// SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(ptr addrspace(4) noundef {{%[0-9]+}}, ptr addrspace(4) noundef {{%[^,]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 2)
// ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(ptr noundef {{%[0-9]+}}, ptr noundef {{%[^,]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 2)
x = __opencl_atomic_compare_exchange_weak(i, &cmp, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_device);
// SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* noundef {{%[0-9]+}}, i8 addrspace(4)* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 3)
// ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8* noundef {{%[0-9]+}}, i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 3)
// SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(ptr addrspace(4) noundef {{%[0-9]+}}, ptr addrspace(4) noundef {{%[^,]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 3)
// ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(ptr noundef {{%[0-9]+}}, ptr noundef {{%[^,]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 3)
x = __opencl_atomic_compare_exchange_weak(i, &cmp, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_all_svm_devices);
#ifdef cl_khr_subgroups
// SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* noundef {{%[0-9]+}}, i8 addrspace(4)* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 4)
// SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(ptr addrspace(4) noundef {{%[0-9]+}}, ptr addrspace(4) noundef {{%[^,]+}}, i32 noundef {{%[0-9]+}}, i32 noundef 5, i32 noundef 5, i32 noundef 4)
x = __opencl_atomic_compare_exchange_weak(i, &cmp, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_sub_group);
#endif
// SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* noundef {{%[0-9]+}}, i8 addrspace(4)* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef %{{.*}}, i32 noundef %{{.*}}, i32 noundef %{{.*}})
// ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8* noundef {{%[0-9]+}}, i8* noundef {{%[0-9]+}}, i32 noundef {{%[0-9]+}}, i32 noundef %{{.*}}, i32 noundef %{{.*}}, i32 noundef %{{.*}})
// SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(ptr addrspace(4) noundef {{%[0-9]+}}, ptr addrspace(4) noundef {{%[^,]+}}, i32 noundef {{%[0-9]+}}, i32 noundef %{{.*}}, i32 noundef %{{.*}}, i32 noundef %{{.*}})
// ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(ptr noundef {{%[0-9]+}}, ptr noundef {{%[^,]+}}, i32 noundef {{%[0-9]+}}, i32 noundef %{{.*}}, i32 noundef %{{.*}}, i32 noundef %{{.*}})
x = __opencl_atomic_compare_exchange_weak(i, &cmp, 1, order, order, scope);
}

90
clang/test/CodeGenOpenCL/atomic-ops.cl
View File

@ -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 @@ atomic_int j;
void fi1(atomic_int *i) {
// CHECK-LABEL: @fi1
// CHECK: load atomic i32, i32* %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
// CHECK: load atomic i32, ptr %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
int x = __opencl_atomic_load(i, memory_order_seq_cst, memory_scope_work_group);
// CHECK: load atomic i32, i32* %{{[.0-9A-Z_a-z]+}} syncscope("agent") seq_cst, align 4
// CHECK: load atomic i32, ptr %{{[.0-9A-Z_a-z]+}} syncscope("agent") seq_cst, align 4
x = __opencl_atomic_load(i, memory_order_seq_cst, memory_scope_device);
// CHECK: load atomic i32, i32* %{{[.0-9A-Z_a-z]+}} seq_cst, align 4
// CHECK: load atomic i32, ptr %{{[.0-9A-Z_a-z]+}} seq_cst, align 4
x = __opencl_atomic_load(i, memory_order_seq_cst, memory_scope_all_svm_devices);
// CHECK: load atomic i32, i32* %{{[.0-9A-Z_a-z]+}} syncscope("wavefront") seq_cst, align 4
// CHECK: load atomic i32, ptr %{{[.0-9A-Z_a-z]+}} syncscope("wavefront") seq_cst, align 4
x = __opencl_atomic_load(i, memory_order_seq_cst, memory_scope_sub_group);
}
void fi2(atomic_int *i) {
// CHECK-LABEL: @fi2
// CHECK: store atomic i32 %{{[.0-9A-Z_a-z]+}}, i32* %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
// CHECK: store atomic i32 %{{[.0-9A-Z_a-z]+}}, ptr %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
__opencl_atomic_store(i, 1, memory_order_seq_cst, memory_scope_work_group);
}
void test_addr(global atomic_int *ig, private atomic_int *ip, local atomic_int *il) {
// CHECK-LABEL: @test_addr
// CHECK: store atomic i32 %{{[.0-9A-Z_a-z]+}}, i32 addrspace(1)* %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
// CHECK: store atomic i32 %{{[.0-9A-Z_a-z]+}}, ptr addrspace(1) %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
__opencl_atomic_store(ig, 1, memory_order_seq_cst, memory_scope_work_group);
// CHECK: store atomic i32 %{{[.0-9A-Z_a-z]+}}, i32 addrspace(5)* %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
// CHECK: store atomic i32 %{{[.0-9A-Z_a-z]+}}, ptr addrspace(5) %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
__opencl_atomic_store(ip, 1, memory_order_seq_cst, memory_scope_work_group);
// CHECK: store atomic i32 %{{[.0-9A-Z_a-z]+}}, i32 addrspace(3)* %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
// CHECK: store atomic i32 %{{[.0-9A-Z_a-z]+}}, ptr addrspace(3) %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
__opencl_atomic_store(il, 1, memory_order_seq_cst, memory_scope_work_group);
}
void fi3(atomic_int *i, atomic_uint *ui) {
// CHECK-LABEL: @fi3
// CHECK: atomicrmw and i32* %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
// CHECK: atomicrmw and ptr %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
int x = __opencl_atomic_fetch_and(i, 1, memory_order_seq_cst, memory_scope_work_group);
// CHECK: atomicrmw min i32* %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
// CHECK: atomicrmw min ptr %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
x = __opencl_atomic_fetch_min(i, 1, memory_order_seq_cst, memory_scope_work_group);
// CHECK: atomicrmw max i32* %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
// CHECK: atomicrmw max ptr %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
x = __opencl_atomic_fetch_max(i, 1, memory_order_seq_cst, memory_scope_work_group);
// CHECK: atomicrmw umin i32* %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
// CHECK: atomicrmw umin ptr %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
x = __opencl_atomic_fetch_min(ui, 1, memory_order_seq_cst, memory_scope_work_group);
// CHECK: atomicrmw umax i32* %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
// CHECK: atomicrmw umax ptr %{{[.0-9A-Z_a-z]+}}, i32 %{{[.0-9A-Z_a-z]+}} syncscope("workgroup") seq_cst, align 4
x = __opencl_atomic_fetch_max(ui, 1, memory_order_seq_cst, memory_scope_work_group);
}
bool fi4(atomic_int *i) {
// CHECK-LABEL: @fi4(
// CHECK: [[PAIR:%[.0-9A-Z_a-z]+]] = cmpxchg i32* [[PTR:%[.0-9A-Z_a-z]+]], i32 [[EXPECTED:%[.0-9A-Z_a-z]+]], i32 [[DESIRED:%[.0-9A-Z_a-z]+]] syncscope("workgroup-one-as") acquire acquire, align 4
// CHECK: [[PAIR:%[.0-9A-Z_a-z]+]] = cmpxchg ptr [[PTR:%[.0-9A-Z_a-z]+]], i32 [[EXPECTED:%[.0-9A-Z_a-z]+]], i32 [[DESIRED:%[.0-9A-Z_a-z]+]] syncscope("workgroup-one-as") acquire acquire, align 4
// CHECK: [[OLD:%[.0-9A-Z_a-z]+]] = extractvalue { i32, i1 } [[PAIR]], 0
// CHECK: [[CMP:%[.0-9A-Z_a-z]+]] = extractvalue { i32, i1 } [[PAIR]], 1
// CHECK: br i1 [[CMP]], label %[[STORE_EXPECTED:[.0-9A-Z_a-z]+]], label %[[CONTINUE:[.0-9A-Z_a-z]+]]
@ -105,16 +105,16 @@ void fi5(atomic_int *i, int scope) {
// 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 @@ void fi6(atomic_int *i, int order, int scope) {
// 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 @@ void ff2(atomic_float *d) {
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 @@ void atomic_init_foo()
// 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 @@ void generalFailureOrder(atomic_int *ptr, int *ptr2, int success, int fail) {
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);
}

96
clang/test/CodeGenOpenCL/blocks.cl
View File

@ -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 @@ void foo(){
// 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 @@ void foo(){
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

17
clang/test/CodeGenOpenCL/builtins.cl
View File

@ -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 @@ void testBranchingOnAddressSpaceCast(generic long* ptr) {
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]]
}

8
clang/test/CodeGenOpenCL/cast_image.cl
View File

@ -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*);
}

6
clang/test/CodeGenOpenCL/const-str-array-decay.cl
View File

@ -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 @@ kernel void str_array_decy() {
test_func("Test string literal");
}
// CHECK: i8 addrspace(2)* noundef getelementptr inbounds ([20 x i8], [20 x i8] addrspace(2)*
// CHECK-NOT: addrspacecast
// CHECK: call i32 @test_func(ptr addrspace(2) noundef @{{.*}})

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

@ -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 @@ typedef struct test_struct {
kernel void test_single(int_single input, global int* output) {
// CHECK: spir_kernel
// AMDGCN: define{{.*}} amdgpu_kernel void @test_single
// CHECK: struct.int_single* nocapture {{.*}} byval(%struct.int_single)
// CHECK: i32* nocapture noundef writeonly align 4 %output
// CHECK: ptr nocapture {{.*}} byval(%struct.int_single)
// CHECK: ptr nocapture noundef writeonly align 4 %output
output[0] = input.a;
}
kernel void test_pair(int_pair input, global int* output) {
// CHECK: spir_kernel
// AMDGCN: define{{.*}} amdgpu_kernel void @test_pair
// CHECK: struct.int_pair* nocapture {{.*}} byval(%struct.int_pair)
// CHECK: i32* nocapture noundef writeonly align 4 %output
// CHECK: ptr nocapture {{.*}} byval(%struct.int_pair)
// CHECK: ptr nocapture noundef writeonly align 4 %output
output[0] = (int)input.a;
output[1] = (int)input.b;
}
@ -44,8 +44,8 @@ kernel void test_pair(int_pair input, global int* output) {
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 @@ kernel void test_kernel(test_struct input, global int* output) {
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;
}

26
clang/test/CodeGenOpenCL/no-half.cl
View File

@ -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
}

46
clang/test/CodeGenOpenCL/pipe_builtin.cl
View File

@ -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);
}

30
clang/test/CodeGenOpenCL/pipe_types.cl
View File

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

87
clang/test/CodeGenOpenCL/to_addr_builtin.cl
View File

@ -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 @@ void test(void) {
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;

9
clang/test/CodeGenSYCL/field-annotate-addr-space.cpp
View File

@ -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 @@ struct HasField {
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;
}

38
clang/test/CodeGenSYCL/unique_stable_name.cpp
View File

@ -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 @@ template <typename KernelName, typename KernelType>
int main() {
kernel_single_task<class kernel2>(func<Derp>);
// CHECK: call spir_func void @_Z18kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_(i8 addrspace(4)* ()* noundef @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv)
// CHECK: call spir_func void @_Z18kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_(ptr noundef @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv)
auto l1 = []() { return 1; };
auto l2 = [](decltype(l1) *l = nullptr) { return 2; };
kernel_single_task<class kernel3>(l2);
puts(__builtin_sycl_unique_stable_name(decltype(l2)));
// CHECK: call spir_func void @_Z18kernel_single_taskIZ4mainE7kernel3Z4mainEUlPZ4mainEUlvE_E_EvT0_
// CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[LAMBDA_K3_SIZE]], [[LAMBDA_K3_SIZE]] addrspace(1)* @[[LAMBDA_KERNEL3]], i32 0, i32 0) to i8 addrspace(4)*))
// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[LAMBDA_KERNEL3]] to ptr addrspace(4)))
constexpr const char str[] = "lalala";
static_assert(__builtin_strcmp(__builtin_sycl_unique_stable_name(decltype(str)), "_ZTSA7_Kc\0") == 0, "unexpected mangling");
int i = 0;
puts(__builtin_sycl_unique_stable_name(decltype(i++)));
// CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]] addrspace(1)* @[[INT1]], i32 0, i32 0) to i8 addrspace(4)*))
// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[INT1]] to ptr addrspace(4)))
// FIXME: Ensure that j is incremented because VLAs are terrible.
int j = 55;
puts(__builtin_sycl_unique_stable_name(int[++j]));
// CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[STRING_SIZE]], [[STRING_SIZE]] addrspace(1)* @[[STRING]], i32 0, i32 0) to i8 addrspace(4)*))
// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[STRING]] to ptr addrspace(4)))
// CHECK: define internal spir_func void @_Z18kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_
// CHECK: declare spir_func noundef i8 addrspace(4)* @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv
// CHECK: declare spir_func noundef ptr addrspace(4) @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv
// CHECK: define internal spir_func void @_Z18kernel_single_taskIZ4mainE7kernel3Z4mainEUlPZ4mainEUlvE_E_EvT0_
// CHECK: define internal spir_func void @_Z18kernel_single_taskIZ4mainE6kernelZ4mainEUlvE0_EvT0_
kernel_single_task<class kernel>(
[]() {
puts(__builtin_sycl_unique_stable_name(int));
// CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]] addrspace(1)* @[[INT2]], i32 0, i32 0) to i8 addrspace(4)*))
// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[INT2]] to ptr addrspace(4)))
auto x = []() {};
puts(__builtin_sycl_unique_stable_name(decltype(x)));
// CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]] addrspace(1)* @[[LAMBDA_X]], i32 0, i32 0) to i8 addrspace(4)*))
// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[LAMBDA_X]] to ptr addrspace(4)))
DEF_IN_MACRO();
// CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]] addrspace(1)* @[[MACRO_X]], i32 0, i32 0) to i8 addrspace(4)*))
// CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]] addrspace(1)* @[[MACRO_Y]], i32 0, i32 0) to i8 addrspace(4)*))
// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[MACRO_X]] to ptr addrspace(4)))
// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[MACRO_Y]] to ptr addrspace(4)))
MACRO_CALLS_MACRO();
// CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]] addrspace(1)* @[[MACRO_MACRO_X]], i32 0, i32 0) to i8 addrspace(4)*))
// CHECK: call spir_func void @puts(i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]] addrspace(1)* @[[MACRO_MACRO_Y]], i32 0, i32 0) to i8 addrspace(4)*))
// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[MACRO_MACRO_X]] to ptr addrspace(4)))
// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[MACRO_MACRO_Y]] to ptr addrspace(4)))
template_param<int>();
// CHECK: call spir_func void @_Z14template_paramIiEvv
@ -138,22 +138,22 @@ int main() {
}
// 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)))

12
clang/test/PCH/arc-blocks.mm
View File

@ -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 @@ namespace test_block_retain {
#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 @@ namespace test_block_retain {
}
// 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);

8
clang/test/Profile/objc-general.m
View File

@ -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 @@ struct NSFastEnumerationState;
@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 @@ struct NSFastEnumerationState;
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

6
clang/test/Rewriter/rewrite-super-message.mm
View File

@ -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 @@ void *sel_registerName(const char *);
}
@end
// CHECK: call %struct.objc_class* @class_getSuperclass
// CHECK: call ptr @class_getSuperclass
@class NSZone;

23
clang/test/SemaObjC/arc-objc-lifetime-conflict.m
View File

@ -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