diff options
| author | Peter Collingbourne <peter@pcc.me.uk> | 2012-05-20 21:08:35 +0000 |
|---|---|---|
| committer | Peter Collingbourne <peter@pcc.me.uk> | 2012-05-20 21:08:35 +0000 |
| commit | f44bdf9c5fc347f1639d50c49de2b836f28017e2 (patch) | |
| tree | a21092fbd172f3ac505a42a8bc578d9daeee9b77 /clang/test/CodeGenCUDA | |
| parent | 853a8a982bb9b3220e3625b725b15faa8ba91d99 (diff) | |
| download | bcm5719-llvm-f44bdf9c5fc347f1639d50c49de2b836f28017e2.tar.gz bcm5719-llvm-f44bdf9c5fc347f1639d50c49de2b836f28017e2.zip | |
CUDA: add CodeGen support for global variable address spaces.
Because in CUDA types do not have associated address spaces,
globals are declared in their "native" address space, and accessed
by bitcasting the pointer to address space 0. This relies on address
space 0 being a unified address space.
llvm-svn: 157167
Diffstat (limited to 'clang/test/CodeGenCUDA')
| -rw-r--r-- | clang/test/CodeGenCUDA/address-spaces.cu | 24 |
1 files changed, 24 insertions, 0 deletions
diff --git a/clang/test/CodeGenCUDA/address-spaces.cu b/clang/test/CodeGenCUDA/address-spaces.cu new file mode 100644 index 00000000000..2da61ec95a6 --- /dev/null +++ b/clang/test/CodeGenCUDA/address-spaces.cu @@ -0,0 +1,24 @@ +// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple ptx32-unknown-unknown | FileCheck %s + +#include "../SemaCUDA/cuda.h" + +// CHECK: @i = global +__device__ int i; + +// CHECK: @j = addrspace(1) global +__constant__ int j; + +// CHECK: @k = addrspace(4) global +__shared__ int k; + +__device__ void foo() { + // CHECK: load i32* @i + i++; + + // CHECK: load i32* bitcast (i32 addrspace(1)* @j to i32*) + j++; + + // CHECK: load i32* bitcast (i32 addrspace(4)* @k to i32*) + k++; +} + |
