1 // RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx-unknown-unknown | FileCheck %s
2 // RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple amdgcn | FileCheck %s
3
4 // Verifies Clang emits correct address spaces and addrspacecast instructions
5 // for CUDA code.
6
7 #include "Inputs/cuda.h"
8
9 // CHECK: @i = addrspace(1) externally_initialized global
10 __device__ int i;
11
12 // CHECK: @j = addrspace(4) externally_initialized global
13 __constant__ int j;
14
15 // CHECK: @k = addrspace(3) global
16 __shared__ int k;
17
18 struct MyStruct {
19 int data1;
20 int data2;
21 };
22
23 // CHECK: @_ZZ5func0vE1a = internal addrspace(3) global %struct.MyStruct undef
24 // CHECK: @_ZZ5func1vE1a = internal addrspace(3) global float undef
25 // CHECK: @_ZZ5func2vE1a = internal addrspace(3) global [256 x float] undef
26 // CHECK: @_ZZ5func3vE1a = internal addrspace(3) global float undef
27 // CHECK: @_ZZ5func4vE1a = internal addrspace(3) global float undef
28 // CHECK: @b = addrspace(3) global float undef
29
foo()30 __device__ void foo() {
31 // CHECK: load i32, i32* addrspacecast (i32 addrspace(1)* @i to i32*)
32 i++;
33
34 // CHECK: load i32, i32* addrspacecast (i32 addrspace(4)* @j to i32*)
35 j++;
36
37 // CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @k to i32*)
38 k++;
39
40 __shared__ int lk;
41 // CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @_ZZ3foovE2lk to i32*)
42 lk++;
43 }
44
func0()45 __device__ void func0() {
46 __shared__ MyStruct a;
47 MyStruct *ap = &a; // composite type
48 ap->data1 = 1;
49 ap->data2 = 2;
50 }
51 // CHECK: define void @_Z5func0v()
52 // CHECK: store %struct.MyStruct* addrspacecast (%struct.MyStruct addrspace(3)* @_ZZ5func0vE1a to %struct.MyStruct*), %struct.MyStruct** %{{.*}}
53
callee(float * ap)54 __device__ void callee(float *ap) {
55 *ap = 1.0f;
56 }
57
func1()58 __device__ void func1() {
59 __shared__ float a;
60 callee(&a); // implicit cast from parameters
61 }
62 // CHECK: define void @_Z5func1v()
63 // CHECK: call void @_Z6calleePf(float* addrspacecast (float addrspace(3)* @_ZZ5func1vE1a to float*))
64
func2()65 __device__ void func2() {
66 __shared__ float a[256];
67 float *ap = &a[128]; // implicit cast from a decayed array
68 *ap = 1.0f;
69 }
70 // CHECK: define void @_Z5func2v()
71 // CHECK: store float* getelementptr inbounds ([256 x float], [256 x float]* addrspacecast ([256 x float] addrspace(3)* @_ZZ5func2vE1a to [256 x float]*), i{{32|64}} 0, i{{32|64}} 128), float** %{{.*}}
72
func3()73 __device__ void func3() {
74 __shared__ float a;
75 float *ap = reinterpret_cast<float *>(&a); // explicit cast
76 *ap = 1.0f;
77 }
78 // CHECK: define void @_Z5func3v()
79 // CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func3vE1a to float*), float** %{{.*}}
80
func4()81 __device__ void func4() {
82 __shared__ float a;
83 float *ap = (float *)&a; // explicit c-style cast
84 *ap = 1.0f;
85 }
86 // CHECK: define void @_Z5func4v()
87 // CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func4vE1a to float*), float** %{{.*}}
88
89 __shared__ float b;
90
func5()91 __device__ float *func5() {
92 return &b; // implicit cast from a return value
93 }
94 // CHECK: define float* @_Z5func5v()
95 // CHECK: ret float* addrspacecast (float addrspace(3)* @b to float*)
96