From f44bdf9c5fc347f1639d50c49de2b836f28017e2 Mon Sep 17 00:00:00 2001 From: Peter Collingbourne Date: Sun, 20 May 2012 21:08:35 +0000 Subject: 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 --- clang/test/CodeGenCUDA/address-spaces.cu | 24 ++++++++++++++++++++++++ 1 file changed, 24 insertions(+) create mode 100644 clang/test/CodeGenCUDA/address-spaces.cu (limited to 'clang/test/CodeGenCUDA') 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++; +} + -- cgit v1.2.3