1 // REQUIRES: x86-registered-target
2 // REQUIRES: amdgpu-registered-target
4 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck --check-prefixes=COMMON,CHECK %s
5 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -disable-O0-optnone -o - | opt -S -O2 | FileCheck %s --check-prefixes=COMMON,OPT
6 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -x hip %s -o - | FileCheck -check-prefix=HOST %s
8 #include "Inputs/cuda.h"
10 // Coerced struct from `struct S` without all generic pointers lowered into
13 // On the host-side compilation, generic pointer won't be coerced.
14 // HOST-NOT: %struct.S.coerce
15 // HOST-NOT: %struct.T.coerce
17 // HOST: define{{.*}} void @_Z22__device_stub__kernel1Pi(ptr noundef %x)
18 // COMMON-LABEL: define{{.*}} amdgpu_kernel void @_Z7kernel1Pi(ptr addrspace(1){{.*}} %x.coerce)
19 // CHECK: ={{.*}} addrspacecast ptr addrspace(1) %{{.*}} to ptr
20 // CHECK-NOT: ={{.*}} addrspacecast ptr addrspace(1) %{{.*}} to ptr
21 // OPT: [[VAL:%.*]] = load i32, ptr addrspace(1) %x.coerce, align 4, !amdgpu.noclobber ![[MD:[0-9]+]]
22 // OPT: [[INC:%.*]] = add nsw i32 [[VAL]], 1
23 // OPT: store i32 [[INC]], ptr addrspace(1) %x.coerce, align 4
25 __global__ void kernel1(int *x) {
29 // HOST: define{{.*}} void @_Z22__device_stub__kernel2Ri(ptr noundef nonnull align 4 dereferenceable(4) %x)
30 // COMMON-LABEL: define{{.*}} amdgpu_kernel void @_Z7kernel2Ri(ptr addrspace(1){{.*}} nonnull align 4 dereferenceable(4) %x.coerce)
31 // CHECK: ={{.*}} addrspacecast ptr addrspace(1) %{{.*}} to ptr
32 // CHECK-NOT: ={{.*}} addrspacecast ptr addrspace(1) %{{.*}} to ptr
33 // OPT: [[VAL:%.*]] = load i32, ptr addrspace(1) %x.coerce, align 4, !amdgpu.noclobber ![[MD]]
34 // OPT: [[INC:%.*]] = add nsw i32 [[VAL]], 1
35 // OPT: store i32 [[INC]], ptr addrspace(1) %x.coerce, align 4
37 __global__ void kernel2(int &x) {
41 // HOST: define{{.*}} void @_Z22__device_stub__kernel3PU3AS2iPU3AS1i(ptr addrspace(2) noundef %x, ptr addrspace(1) noundef %y)
42 // CHECK-LABEL: define{{.*}} amdgpu_kernel void @_Z7kernel3PU3AS2iPU3AS1i(ptr addrspace(2){{.*}} %x, ptr addrspace(1){{.*}} %y)
43 // CHECK-NOT: ={{.*}} addrspacecast ptr addrspace(1) %{{.*}} to ptr
44 __global__ void kernel3(__attribute__((address_space(2))) int *x,
45 __attribute__((address_space(1))) int *y) {
49 // COMMON-LABEL: define{{.*}} void @_Z4funcPi(ptr{{.*}} %x)
50 // CHECK-NOT: ={{.*}} addrspacecast ptr addrspace(1) %{{.*}} to ptr
51 __device__ void func(int *x) {
59 // `by-val` struct is passed by-indirect-alias (a mix of by-ref and indirect
60 // by-val). However, the enhanced address inferring pass should be able to
61 // assume they are global pointers.
63 // HOST: define{{.*}} void @_Z22__device_stub__kernel41S(ptr %s.coerce0, ptr %s.coerce1)
64 // COMMON-LABEL: define{{.*}} amdgpu_kernel void @_Z7kernel41S(ptr addrspace(4){{.*}} byref(%struct.S) align 8 %0)
65 // OPT: [[P0:%.*]] = load ptr, ptr addrspace(4) %0, align 8
66 // OPT: [[G0:%.*]] ={{.*}} addrspacecast ptr [[P0]] to ptr addrspace(1)
67 // OPT: [[R1:%.*]] = getelementptr inbounds i8, ptr addrspace(4) %0, i64 8
68 // OPT: [[P1:%.*]] = load ptr, ptr addrspace(4) [[R1]], align 8
69 // OPT: [[G1:%.*]] ={{.*}} addrspacecast ptr [[P1]] to ptr addrspace(1)
70 // OPT: [[V0:%.*]] = load i32, ptr addrspace(1) [[G0]], align 4, !amdgpu.noclobber ![[MD]]
71 // OPT: [[INC:%.*]] = add nsw i32 [[V0]], 1
72 // OPT: store i32 [[INC]], ptr addrspace(1) [[G0]], align 4
73 // OPT: [[V1:%.*]] = load float, ptr addrspace(1) [[G1]], align 4
74 // OPT: [[ADD:%.*]] = fadd contract float [[V1]], 1.000000e+00
75 // OPT: store float [[ADD]], ptr addrspace(1) [[G1]], align 4
77 __global__ void kernel4(struct S s) {
82 // If a pointer to struct is passed, only the pointer itself is coerced into the global one.
83 // HOST: define{{.*}} void @_Z22__device_stub__kernel5P1S(ptr noundef %s)
84 // COMMON-LABEL: define{{.*}} amdgpu_kernel void @_Z7kernel5P1S(ptr addrspace(1){{.*}} %s.coerce)
85 __global__ void kernel5(struct S *s) {
93 // `by-val` array is passed by-indirect-alias (a mix of by-ref and indirect
94 // by-val). However, the enhanced address inferring pass should be able to
95 // assume they are global pointers.
97 // HOST: define{{.*}} void @_Z22__device_stub__kernel61T(ptr %t.coerce0, ptr %t.coerce1)
98 // COMMON-LABEL: define{{.*}} amdgpu_kernel void @_Z7kernel61T(ptr addrspace(4){{.*}} byref(%struct.T) align 8 %0)
99 // OPT: [[P0:%.*]] = load ptr, ptr addrspace(4) %0, align 8
100 // OPT: [[G0:%.*]] ={{.*}} addrspacecast ptr [[P0]] to ptr addrspace(1)
101 // OPT: [[R1:%.*]] = getelementptr inbounds i8, ptr addrspace(4) %0, i64 8
102 // OPT: [[P1:%.*]] = load ptr, ptr addrspace(4) [[R1]], align 8
103 // OPT: [[G1:%.*]] ={{.*}} addrspacecast ptr [[P1]] to ptr addrspace(1)
104 // OPT: [[V0:%.*]] = load float, ptr addrspace(1) [[G0]], align 4, !amdgpu.noclobber ![[MD]]
105 // OPT: [[ADD0:%.*]] = fadd contract float [[V0]], 1.000000e+00
106 // OPT: store float [[ADD0]], ptr addrspace(1) [[G0]], align 4
107 // OPT: [[V1:%.*]] = load float, ptr addrspace(1) [[G1]], align 4
108 // OPT: [[ADD1:%.*]] = fadd contract float [[V1]], 2.000000e+00
109 // OPT: store float [[ADD1]], ptr addrspace(1) [[G1]], align 4
111 __global__ void kernel6(struct T t) {
116 // Check that coerced pointers retain the noalias attribute when qualified with __restrict.
117 // HOST: define{{.*}} void @_Z22__device_stub__kernel7Pi(ptr noalias noundef %x)
118 // COMMON-LABEL: define{{.*}} amdgpu_kernel void @_Z7kernel7Pi(ptr addrspace(1) noalias{{.*}} %x.coerce)
119 __global__ void kernel7(int *__restrict x) {
123 // Single element struct.
127 // HOST: define{{.*}} void @_Z22__device_stub__kernel82SS(ptr %a.coerce)
128 // COMMON-LABEL: define{{.*}} amdgpu_kernel void @_Z7kernel82SS(ptr addrspace(1){{.*}} %a.coerce)
129 // CHECK: ={{.*}} addrspacecast ptr addrspace(1) %{{.*}} to ptr
130 // CHECK-NOT: ={{.*}} addrspacecast ptr addrspace(1) %{{.*}} to ptr
131 // OPT: [[VAL:%.*]] = load float, ptr addrspace(1) %a.coerce, align 4, !amdgpu.noclobber ![[MD]]
132 // OPT: [[INC:%.*]] = fadd contract float [[VAL]], 3.000000e+00
133 // OPT: store float [[INC]], ptr addrspace(1) %a.coerce, align 4
135 __global__ void kernel8(struct SS a) {