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