• Home
  • Line#
  • Scopes#
  • Navigate#
  • Raw
  • Download
1 // RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx-unknown-unknown | FileCheck %s
2 
3 #include "../SemaCUDA/cuda.h"
4 
5 // CHECK: @i = addrspace(1) global
6 __device__ int i;
7 
8 // CHECK: @j = addrspace(4) global
9 __constant__ int j;
10 
11 // CHECK: @k = addrspace(3) global
12 __shared__ int k;
13 
foo()14 __device__ void foo() {
15   // CHECK: load i32* bitcast (i32 addrspace(1)* @i to i32*)
16   i++;
17 
18   // CHECK: load i32* bitcast (i32 addrspace(4)* @j to i32*)
19   j++;
20 
21   // CHECK: load i32* bitcast (i32 addrspace(3)* @k to i32*)
22   k++;
23 
24   static int li;
25   // CHECK: load i32 addrspace(1)* @_ZZ3foovE2li
26   li++;
27 
28   __constant__ int lj;
29   // CHECK: load i32 addrspace(4)* @_ZZ3foovE2lj
30   lj++;
31 
32   __shared__ int lk;
33   // CHECK: load i32 addrspace(3)* @_ZZ3foovE2lk
34   lk++;
35 }
36 
37