ptx: add state spaces

git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@122638 91177308-0d34-0410-b5e6-96231b3b80d8

1 files changed, 34 insertions, 1 deletions

diff --git a/test/CodeGen/PTX/ld.ll b/test/CodeGen/PTX/ld.ll
index 5ae82a0385..baafbc2d3d 100644
--- a/test/CodeGen/PTX/ld.ll
+++ b/test/CodeGen/PTX/ld.ll
@@ -3,6 +3,15 @@
 ;CHECK: .extern .global .s32 array[];
 @array = external global [10 x i32]
 
+;CHECK: .extern .const .s32 array_constant[];
+@array_constant = external addrspace(1) constant [10 x i32]
+
+;CHECK: .extern .local .s32 array_local[];
+@array_local = external addrspace(2) global [10 x i32]
+
+;CHECK: .extern .shared .s32 array_shared[];
+@array_shared = external addrspace(4) global [10 x i32]
+
 define ptx_device i32 @t1(i32* %p) {
 entry:
 ;CHECK: ld.global.s32 r0, [r1];
@@ -27,7 +36,7 @@ entry:
   ret i32 %x
 }
 
-define ptx_device i32 @t4() {
+define ptx_device i32 @t4_global() {
 entry:
 ;CHECK: ld.global.s32 r0, [array];
   %i = getelementptr [10 x i32]* @array, i32 0, i32 0
@@ -35,6 +44,30 @@ entry:
   ret i32 %x
 }
 
+define ptx_device i32 @t4_const() {
+entry:
+;CHECK: ld.const.s32 r0, [array_constant];
+  %i = getelementptr [10 x i32] addrspace(1)* @array_constant, i32 0, i32 0
+  %x = load i32 addrspace(1)* %i
+  ret i32 %x
+}
+
+define ptx_device i32 @t4_local() {
+entry:
+;CHECK: ld.local.s32 r0, [array_local];
+  %i = getelementptr [10 x i32] addrspace(2)* @array_local, i32 0, i32 0
+  %x = load i32 addrspace(2)* %i
+  ret i32 %x
+}
+
+define ptx_device i32 @t4_shared() {
+entry:
+;CHECK: ld.shared.s32 r0, [array_shared];
+  %i = getelementptr [10 x i32] addrspace(4)* @array_shared, i32 0, i32 0
+  %x = load i32 addrspace(4)* %i
+  ret i32 %x
+}
+
 define ptx_device i32 @t5() {
 entry:
 ;CHECK: ld.global.s32 r0, [array+4];