Skip to content

Commit f1885fd

Browse files
AlexMacleanAlexMaclean
committed
pre-commit tests
1 parent 414590b commit f1885fd

File tree

2 files changed

+86
-15
lines changed

2 files changed

+86
-15
lines changed

‎llvm/test/CodeGen/NVPTX/bug26185-2.ll

Lines changed: 21 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,4 @@
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
12
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_35 -verify-machineinstrs | FileCheck %s
23
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_35 | %ptxas-verify %}
34

@@ -10,14 +11,31 @@
1011
target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
1112
target triple = "nvptx64-nvidia-cuda"
1213

13-
; CHECK-LABEL: spam
1414
define ptx_kernel void @spam(ptr addrspace(1) noalias nocapture readonly %arg, ptr addrspace(1) noalias nocapture %arg1, i64 %arg2, i64 %arg3) #0 {
15+
; CHECK-LABEL: spam(
16+
; CHECK: .maxntid 1, 1, 1
17+
; CHECK-NEXT: {
18+
; CHECK-NEXT: .reg .b16 %rs<2>;
19+
; CHECK-NEXT: .reg .b32 %r<2>;
20+
; CHECK-NEXT: .reg .b64 %rd<9>;
21+
; CHECK-EMPTY:
22+
; CHECK-NEXT: // %bb.0: // %bb
23+
; CHECK-NEXT: ld.param.b64 %rd1, [spam_param_0];
24+
; CHECK-NEXT: ld.param.b64 %rd2, [spam_param_3];
25+
; CHECK-NEXT: shl.b64 %rd3, %rd2, 1;
26+
; CHECK-NEXT: add.s64 %rd4, %rd1, %rd3;
27+
; CHECK-NEXT: ld.param.b64 %rd5, [spam_param_1];
28+
; CHECK-NEXT: ld.global.nc.b16 %rs1, [%rd4+16];
29+
; CHECK-NEXT: cvt.s32.s16 %r1, %rs1;
30+
; CHECK-NEXT: mul.wide.s32 %rd6, %r1, %r1;
31+
; CHECK-NEXT: ld.global.b64 %rd7, [%rd5];
32+
; CHECK-NEXT: add.s64 %rd8, %rd6, %rd7;
33+
; CHECK-NEXT: st.global.b64 [%rd5], %rd8;
34+
; CHECK-NEXT: ret;
1535
bb:
1636
%tmp5 = add nsw i64 %arg3, 8
1737
%tmp6 = getelementptr i16, ptr addrspace(1) %arg, i64 %tmp5
18-
; CHECK: ld.global.nc.b16
1938
%tmp7 = load i16, ptr addrspace(1) %tmp6, align 2
20-
; CHECK: cvt.s32.s16
2139
%tmp8 = sext i16 %tmp7 to i64
2240
%tmp9 = mul nsw i64 %tmp8, %tmp8
2341
%tmp10 = load i64, ptr addrspace(1) %arg1, align 8

‎llvm/test/CodeGen/NVPTX/bug26185.ll

Lines changed: 65 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,4 @@
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
12
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_35 -verify-machineinstrs | FileCheck %s
23
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_35 | %ptxas-verify %}
34

@@ -7,45 +8,97 @@
78
target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
89
target triple = "nvptx64-unknown-unknown"
910

10-
; CHECK-LABEL: ex_zext
1111
define ptx_kernel void @ex_zext(ptr noalias readonly %data, ptr %res) {
12+
; CHECK-LABEL: ex_zext(
13+
; CHECK: {
14+
; CHECK-NEXT: .reg .b16 %rs<2>;
15+
; CHECK-NEXT: .reg .b32 %r<2>;
16+
; CHECK-NEXT: .reg .b64 %rd<5>;
17+
; CHECK-EMPTY:
18+
; CHECK-NEXT: // %bb.0: // %entry
19+
; CHECK-NEXT: ld.param.b64 %rd1, [ex_zext_param_0];
20+
; CHECK-NEXT: cvta.to.global.u64 %rd2, %rd1;
21+
; CHECK-NEXT: ld.param.b64 %rd3, [ex_zext_param_1];
22+
; CHECK-NEXT: cvta.to.global.u64 %rd4, %rd3;
23+
; CHECK-NEXT: ld.global.nc.b8 %rs1, [%rd2];
24+
; CHECK-NEXT: cvt.u32.u8 %r1, %rs1;
25+
; CHECK-NEXT: st.global.b32 [%rd4], %r1;
26+
; CHECK-NEXT: ret;
1227
entry:
13-
; CHECK: ld.global.nc.b8
1428
%val = load i8, ptr %data
15-
; CHECK: cvt.u32.u8
1629
%valext = zext i8 %val to i32
1730
store i32 %valext, ptr %res
1831
ret void
1932
}
2033

21-
; CHECK-LABEL: ex_sext
2234
define ptx_kernel void @ex_sext(ptr noalias readonly %data, ptr %res) {
35+
; CHECK-LABEL: ex_sext(
36+
; CHECK: {
37+
; CHECK-NEXT: .reg .b16 %rs<2>;
38+
; CHECK-NEXT: .reg .b32 %r<2>;
39+
; CHECK-NEXT: .reg .b64 %rd<5>;
40+
; CHECK-EMPTY:
41+
; CHECK-NEXT: // %bb.0: // %entry
42+
; CHECK-NEXT: ld.param.b64 %rd1, [ex_sext_param_0];
43+
; CHECK-NEXT: cvta.to.global.u64 %rd2, %rd1;
44+
; CHECK-NEXT: ld.param.b64 %rd3, [ex_sext_param_1];
45+
; CHECK-NEXT: cvta.to.global.u64 %rd4, %rd3;
46+
; CHECK-NEXT: ld.global.nc.b8 %rs1, [%rd2];
47+
; CHECK-NEXT: cvt.s32.s8 %r1, %rs1;
48+
; CHECK-NEXT: st.global.b32 [%rd4], %r1;
49+
; CHECK-NEXT: ret;
2350
entry:
24-
; CHECK: ld.global.nc.b8
2551
%val = load i8, ptr %data
26-
; CHECK: cvt.s32.s8
2752
%valext = sext i8 %val to i32
2853
store i32 %valext, ptr %res
2954
ret void
3055
}
3156

32-
; CHECK-LABEL: ex_zext_v2
3357
define ptx_kernel void @ex_zext_v2(ptr noalias readonly %data, ptr %res) {
58+
; CHECK-LABEL: ex_zext_v2(
59+
; CHECK: {
60+
; CHECK-NEXT: .reg .b16 %rs<3>;
61+
; CHECK-NEXT: .reg .b32 %r<3>;
62+
; CHECK-NEXT: .reg .b64 %rd<5>;
63+
; CHECK-EMPTY:
64+
; CHECK-NEXT: // %bb.0: // %entry
65+
; CHECK-NEXT: ld.param.b64 %rd1, [ex_zext_v2_param_0];
66+
; CHECK-NEXT: cvta.to.global.u64 %rd2, %rd1;
67+
; CHECK-NEXT: ld.param.b64 %rd3, [ex_zext_v2_param_1];
68+
; CHECK-NEXT: cvta.to.global.u64 %rd4, %rd3;
69+
; CHECK-NEXT: ld.global.nc.v2.b8 {%rs1, %rs2}, [%rd2];
70+
; CHECK-NEXT: cvt.u32.u16 %r1, %rs2;
71+
; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
72+
; CHECK-NEXT: st.global.v2.b32 [%rd4], {%r2, %r1};
73+
; CHECK-NEXT: ret;
3474
entry:
35-
; CHECK: ld.global.nc.v2.b8
3675
%val = load <2 x i8>, ptr %data
37-
; CHECK: cvt.u32.u16
3876
%valext = zext <2 x i8> %val to <2 x i32>
3977
store <2 x i32> %valext, ptr %res
4078
ret void
4179
}
4280

43-
; CHECK-LABEL: ex_sext_v2
4481
define ptx_kernel void @ex_sext_v2(ptr noalias readonly %data, ptr %res) {
82+
; CHECK-LABEL: ex_sext_v2(
83+
; CHECK: {
84+
; CHECK-NEXT: .reg .b16 %rs<3>;
85+
; CHECK-NEXT: .reg .b32 %r<5>;
86+
; CHECK-NEXT: .reg .b64 %rd<5>;
87+
; CHECK-EMPTY:
88+
; CHECK-NEXT: // %bb.0: // %entry
89+
; CHECK-NEXT: ld.param.b64 %rd1, [ex_sext_v2_param_0];
90+
; CHECK-NEXT: cvta.to.global.u64 %rd2, %rd1;
91+
; CHECK-NEXT: ld.param.b64 %rd3, [ex_sext_v2_param_1];
92+
; CHECK-NEXT: cvta.to.global.u64 %rd4, %rd3;
93+
; CHECK-NEXT: ld.global.nc.v2.b8 {%rs1, %rs2}, [%rd2];
94+
; CHECK-NEXT: cvt.u32.u16 %r1, %rs2;
95+
; CHECK-NEXT: cvt.s32.s8 %r2, %r1;
96+
; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
97+
; CHECK-NEXT: cvt.s32.s8 %r4, %r3;
98+
; CHECK-NEXT: st.global.v2.b32 [%rd4], {%r4, %r2};
99+
; CHECK-NEXT: ret;
45100
entry:
46-
; CHECK: ld.global.nc.v2.b8
47101
%val = load <2 x i8>, ptr %data
48-
; CHECK: cvt.s32.s8
49102
%valext = sext <2 x i8> %val to <2 x i32>
50103
store <2 x i32> %valext, ptr %res
51104
ret void

0 commit comments

Comments
 (0)