maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity
llvm-project/clang/test/CodeGenCUDA/address-spaces.cu

37 lines
764 B
Plaintext
Raw Normal View History

Replace PTX back-end with NVPTX back-end in all places where Clang cares NV_CONTRIB llvm-svn: 157403
2012-05-25 01:43:12 +08:00
// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx-unknown-unknown | FileCheck %s
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
2012-05-21 05:08:35 +08:00
#include "../SemaCUDA/cuda.h"
Replace PTX back-end with NVPTX back-end in all places where Clang cares NV_CONTRIB llvm-svn: 157403
2012-05-25 01:43:12 +08:00
// CHECK: @i = addrspace(1) global
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
2012-05-21 05:08:35 +08:00
__device__ int i;
Replace PTX back-end with NVPTX back-end in all places where Clang cares NV_CONTRIB llvm-svn: 157403
2012-05-25 01:43:12 +08:00
// CHECK: @j = addrspace(4) global
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
2012-05-21 05:08:35 +08:00
__constant__ int j;
Replace PTX back-end with NVPTX back-end in all places where Clang cares NV_CONTRIB llvm-svn: 157403
2012-05-25 01:43:12 +08:00
// CHECK: @k = addrspace(3) global
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
2012-05-21 05:08:35 +08:00
__shared__ int k;
__device__ void foo() {
Replace PTX back-end with NVPTX back-end in all places where Clang cares NV_CONTRIB llvm-svn: 157403
2012-05-25 01:43:12 +08:00
// CHECK: load i32* bitcast (i32 addrspace(1)* @i to i32*)
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
2012-05-21 05:08:35 +08:00
i++;
Replace PTX back-end with NVPTX back-end in all places where Clang cares NV_CONTRIB llvm-svn: 157403
2012-05-25 01:43:12 +08:00
// CHECK: load i32* bitcast (i32 addrspace(4)* @j to i32*)
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
2012-05-21 05:08:35 +08:00
j++;
Replace PTX back-end with NVPTX back-end in all places where Clang cares NV_CONTRIB llvm-svn: 157403
2012-05-25 01:43:12 +08:00
// CHECK: load i32* bitcast (i32 addrspace(3)* @k to i32*)
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
2012-05-21 05:08:35 +08:00
k++;
CUDA: give correct address space to globals declared in functions llvm-svn: 162787
2012-08-29 04:37:10 +08:00
static int li;
// CHECK: load i32 addrspace(1)* @_ZZ3foovE2li
li++;
CUDA: give static storage class to __shared__ and __constant__ variables without a storage class within a function, to implement CUDA B.2.5: "__shared__ and __constant__ variables have implied static storage [duration]." llvm-svn: 162788
2012-08-29 04:37:50 +08:00
__constant__ int lj;
// CHECK: load i32 addrspace(4)* @_ZZ3foovE2lj
lj++;
__shared__ int lk;
// CHECK: load i32 addrspace(3)* @_ZZ3foovE2lk
lk++;
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
2012-05-21 05:08:35 +08:00
}