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 14 __device__ void foo() { 15 // CHECK: load i32* addrspacecast (i32 addrspace(1)* @i to i32*) 16 i++; 17 18 // CHECK: load i32* addrspacecast (i32 addrspace(4)* @j to i32*) 19 j++; 20 21 // CHECK: load i32* addrspacecast (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