summaryrefslogtreecommitdiff
path: root/test/CodeGen/PTX
diff options
context:
space:
mode:
Diffstat (limited to 'test/CodeGen/PTX')
-rw-r--r--test/CodeGen/PTX/add.ll15
-rw-r--r--test/CodeGen/PTX/dg.exp5
-rw-r--r--test/CodeGen/PTX/exit.ll14
-rw-r--r--test/CodeGen/PTX/ld.ll78
-rw-r--r--test/CodeGen/PTX/mov.ll13
-rw-r--r--test/CodeGen/PTX/options.ll6
-rw-r--r--test/CodeGen/PTX/ret.ll7
-rw-r--r--test/CodeGen/PTX/shl.ll22
-rw-r--r--test/CodeGen/PTX/shr.ll43
-rw-r--r--test/CodeGen/PTX/st.ll71
-rw-r--r--test/CodeGen/PTX/sub.ll15
11 files changed, 289 insertions, 0 deletions
diff --git a/test/CodeGen/PTX/add.ll b/test/CodeGen/PTX/add.ll
new file mode 100644
index 000000000000..1259d03e96c9
--- /dev/null
+++ b/test/CodeGen/PTX/add.ll
@@ -0,0 +1,15 @@
+; RUN: llc < %s -march=ptx | FileCheck %s
+
+define ptx_device i32 @t1(i32 %x, i32 %y) {
+; CHECK: add.s32 r0, r1, r2;
+ %z = add i32 %x, %y
+; CHECK: ret;
+ ret i32 %z
+}
+
+define ptx_device i32 @t2(i32 %x) {
+; CHECK: add.s32 r0, r1, 1;
+ %z = add i32 %x, 1
+; CHECK: ret;
+ ret i32 %z
+}
diff --git a/test/CodeGen/PTX/dg.exp b/test/CodeGen/PTX/dg.exp
new file mode 100644
index 000000000000..2c304b57741e
--- /dev/null
+++ b/test/CodeGen/PTX/dg.exp
@@ -0,0 +1,5 @@
+load_lib llvm.exp
+
+if { [llvm_supports_target PTX] } {
+ RunLLVMTests [lsort [glob -nocomplain $srcdir/$subdir/*.{ll,c,cpp}]]
+}
diff --git a/test/CodeGen/PTX/exit.ll b/test/CodeGen/PTX/exit.ll
new file mode 100644
index 000000000000..4071babb80ce
--- /dev/null
+++ b/test/CodeGen/PTX/exit.ll
@@ -0,0 +1,14 @@
+; RUN: llc < %s -march=ptx | FileCheck %s
+
+define ptx_kernel void @t1() {
+; CHECK: exit;
+; CHECK-NOT: ret;
+ ret void
+}
+
+define ptx_kernel void @t2(i32* %p, i32 %x) {
+ store i32 %x, i32* %p
+; CHECK: exit;
+; CHECK-NOT: ret;
+ ret void
+}
diff --git a/test/CodeGen/PTX/ld.ll b/test/CodeGen/PTX/ld.ll
new file mode 100644
index 000000000000..836c4d41045a
--- /dev/null
+++ b/test/CodeGen/PTX/ld.ll
@@ -0,0 +1,78 @@
+; RUN: llc < %s -march=ptx | FileCheck %s
+
+;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];
+ %x = load i32* %p
+ ret i32 %x
+}
+
+define ptx_device i32 @t2(i32* %p) {
+entry:
+;CHECK: ld.global.s32 r0, [r1+4];
+ %i = getelementptr i32* %p, i32 1
+ %x = load i32* %i
+ ret i32 %x
+}
+
+define ptx_device i32 @t3(i32* %p, i32 %q) {
+entry:
+;CHECK: shl.b32 r0, r2, 2;
+;CHECK: add.s32 r0, r1, r0;
+;CHECK: ld.global.s32 r0, [r0];
+ %i = getelementptr i32* %p, i32 %q
+ %x = load i32* %i
+ ret i32 %x
+}
+
+define ptx_device i32 @t4_global() {
+entry:
+;CHECK: ld.global.s32 r0, [array];
+ %i = getelementptr [10 x i32]* @array, i32 0, i32 0
+ %x = load i32* %i
+ 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];
+ %i = getelementptr [10 x i32]* @array, i32 0, i32 1
+ %x = load i32* %i
+ ret i32 %x
+}
diff --git a/test/CodeGen/PTX/mov.ll b/test/CodeGen/PTX/mov.ll
new file mode 100644
index 000000000000..c365e9beb897
--- /dev/null
+++ b/test/CodeGen/PTX/mov.ll
@@ -0,0 +1,13 @@
+; RUN: llc < %s -march=ptx | FileCheck %s
+
+define ptx_device i32 @t1() {
+; CHECK: mov.s32 r0, 0;
+; CHECK: ret;
+ ret i32 0
+}
+
+define ptx_device i32 @t2(i32 %x) {
+; CHECK: mov.s32 r0, r1;
+; CHECK: ret;
+ ret i32 %x
+}
diff --git a/test/CodeGen/PTX/options.ll b/test/CodeGen/PTX/options.ll
new file mode 100644
index 000000000000..a14d5c9c27ba
--- /dev/null
+++ b/test/CodeGen/PTX/options.ll
@@ -0,0 +1,6 @@
+; RUN: llc < %s -march=ptx -ptx-version=2.0 | grep ".version 2.0"
+; RUN: llc < %s -march=ptx -ptx-target=sm_20 | grep ".target sm_20"
+
+define ptx_device void @t1() {
+ ret void
+}
diff --git a/test/CodeGen/PTX/ret.ll b/test/CodeGen/PTX/ret.ll
new file mode 100644
index 000000000000..d5037f25fd36
--- /dev/null
+++ b/test/CodeGen/PTX/ret.ll
@@ -0,0 +1,7 @@
+; RUN: llc < %s -march=ptx | FileCheck %s
+
+define ptx_device void @t1() {
+; CHECK: ret;
+; CHECK-NOT: exit;
+ ret void
+}
diff --git a/test/CodeGen/PTX/shl.ll b/test/CodeGen/PTX/shl.ll
new file mode 100644
index 000000000000..b564b43ab932
--- /dev/null
+++ b/test/CodeGen/PTX/shl.ll
@@ -0,0 +1,22 @@
+; RUN: llc < %s -march=ptx | FileCheck %s
+
+define ptx_device i32 @t1(i32 %x, i32 %y) {
+; CHECK: shl.b32 r0, r1, r2
+ %z = shl i32 %x, %y
+; CHECK: ret;
+ ret i32 %z
+}
+
+define ptx_device i32 @t2(i32 %x) {
+; CHECK: shl.b32 r0, r1, 3
+ %z = shl i32 %x, 3
+; CHECK: ret;
+ ret i32 %z
+}
+
+define ptx_device i32 @t3(i32 %x) {
+; CHECK: shl.b32 r0, 3, r1
+ %z = shl i32 3, %x
+; CHECK: ret;
+ ret i32 %z
+}
diff --git a/test/CodeGen/PTX/shr.ll b/test/CodeGen/PTX/shr.ll
new file mode 100644
index 000000000000..3f8ade862b75
--- /dev/null
+++ b/test/CodeGen/PTX/shr.ll
@@ -0,0 +1,43 @@
+; RUN: llc < %s -march=ptx | FileCheck %s
+
+define ptx_device i32 @t1(i32 %x, i32 %y) {
+; CHECK: shr.u32 r0, r1, r2
+ %z = lshr i32 %x, %y
+; CHECK: ret;
+ ret i32 %z
+}
+
+define ptx_device i32 @t2(i32 %x) {
+; CHECK: shr.u32 r0, r1, 3
+ %z = lshr i32 %x, 3
+; CHECK: ret;
+ ret i32 %z
+}
+
+define ptx_device i32 @t3(i32 %x) {
+; CHECK: shr.u32 r0, 3, r1
+ %z = lshr i32 3, %x
+; CHECK: ret;
+ ret i32 %z
+}
+
+define ptx_device i32 @t4(i32 %x, i32 %y) {
+; CHECK: shr.s32 r0, r1, r2
+ %z = ashr i32 %x, %y
+; CHECK: ret;
+ ret i32 %z
+}
+
+define ptx_device i32 @t5(i32 %x) {
+; CHECK: shr.s32 r0, r1, 3
+ %z = ashr i32 %x, 3
+; CHECK: ret;
+ ret i32 %z
+}
+
+define ptx_device i32 @t6(i32 %x) {
+; CHECK: shr.s32 r0, -3, r1
+ %z = ashr i32 -3, %x
+; CHECK: ret;
+ ret i32 %z
+}
diff --git a/test/CodeGen/PTX/st.ll b/test/CodeGen/PTX/st.ll
new file mode 100644
index 000000000000..2cbacb9ee59c
--- /dev/null
+++ b/test/CodeGen/PTX/st.ll
@@ -0,0 +1,71 @@
+; RUN: llc < %s -march=ptx | FileCheck %s
+
+;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 void @t1(i32* %p, i32 %x) {
+entry:
+;CHECK: st.global.s32 [r1], r2;
+ store i32 %x, i32* %p
+ ret void
+}
+
+define ptx_device void @t2(i32* %p, i32 %x) {
+entry:
+;CHECK: st.global.s32 [r1+4], r2;
+ %i = getelementptr i32* %p, i32 1
+ store i32 %x, i32* %i
+ ret void
+}
+
+define ptx_device void @t3(i32* %p, i32 %q, i32 %x) {
+;CHECK: .reg .s32 r0;
+entry:
+;CHECK: shl.b32 r0, r2, 2;
+;CHECK: add.s32 r0, r1, r0;
+;CHECK: st.global.s32 [r0], r3;
+ %i = getelementptr i32* %p, i32 %q
+ store i32 %x, i32* %i
+ ret void
+}
+
+define ptx_device void @t4_global(i32 %x) {
+entry:
+;CHECK: st.global.s32 [array], r1;
+ %i = getelementptr [10 x i32]* @array, i32 0, i32 0
+ store i32 %x, i32* %i
+ ret void
+}
+
+define ptx_device void @t4_local(i32 %x) {
+entry:
+;CHECK: st.local.s32 [array_local], r1;
+ %i = getelementptr [10 x i32] addrspace(2)* @array_local, i32 0, i32 0
+ store i32 %x, i32 addrspace(2)* %i
+ ret void
+}
+
+define ptx_device void @t4_shared(i32 %x) {
+entry:
+;CHECK: st.shared.s32 [array_shared], r1;
+ %i = getelementptr [10 x i32] addrspace(4)* @array_shared, i32 0, i32 0
+ store i32 %x, i32 addrspace(4)* %i
+ ret void
+}
+
+define ptx_device void @t5(i32 %x) {
+entry:
+;CHECK: st.global.s32 [array+4], r1;
+ %i = getelementptr [10 x i32]* @array, i32 0, i32 1
+ store i32 %x, i32* %i
+ ret void
+}
diff --git a/test/CodeGen/PTX/sub.ll b/test/CodeGen/PTX/sub.ll
new file mode 100644
index 000000000000..aab3fdadad13
--- /dev/null
+++ b/test/CodeGen/PTX/sub.ll
@@ -0,0 +1,15 @@
+; RUN: llc < %s -march=ptx | FileCheck %s
+
+define ptx_device i32 @t1(i32 %x, i32 %y) {
+;CHECK: sub.s32 r0, r1, r2;
+ %z = sub i32 %x, %y
+;CHECK: ret;
+ ret i32 %z
+}
+
+define ptx_device i32 @t2(i32 %x) {
+;CHECK: add.s32 r0, r1, -1;
+ %z = sub i32 %x, 1
+;CHECK: ret;
+ ret i32 %z
+}
generated by cgit v1.2.3 (git 2.25.1) at 2025-10-11 02:36:18 +0000