Skip to content

Commit 0002c17

Browse files
AlexMacleanAlexMaclean
committed
pre-commit tests
1 parent f480e1b commit 0002c17

File tree

2 files changed

+137
-0
lines changed

2 files changed

+137
-0
lines changed

‎llvm/test/CodeGen/NVPTX/i8x4-instructions.ll‎

Lines changed: 52 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2311,4 +2311,56 @@ entry:
23112311
ret void
23122312
}
23132313

2314+
define <4 x float> @test_uitofp_v4i8(<4 x i8> %a) {
2315+
; CHECK-LABEL: test_uitofp_v4i8(
2316+
; CHECK: {
2317+
; CHECK-NEXT: .reg .b16 %rs<5>;
2318+
; CHECK-NEXT: .reg .b32 %r<10>;
2319+
; CHECK-EMPTY:
2320+
; CHECK-NEXT: // %bb.0:
2321+
; CHECK-NEXT: ld.param.b32 %r1, [test_uitofp_v4i8_param_0];
2322+
; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U;
2323+
; CHECK-NEXT: cvt.u16.u32 %rs1, %r2;
2324+
; CHECK-NEXT: cvt.rn.f32.u16 %r3, %rs1;
2325+
; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7772U;
2326+
; CHECK-NEXT: cvt.u16.u32 %rs2, %r4;
2327+
; CHECK-NEXT: cvt.rn.f32.u16 %r5, %rs2;
2328+
; CHECK-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;
2329+
; CHECK-NEXT: cvt.u16.u32 %rs3, %r6;
2330+
; CHECK-NEXT: cvt.rn.f32.u16 %r7, %rs3;
2331+
; CHECK-NEXT: prmt.b32 %r8, %r1, 0, 0x7770U;
2332+
; CHECK-NEXT: cvt.u16.u32 %rs4, %r8;
2333+
; CHECK-NEXT: cvt.rn.f32.u16 %r9, %rs4;
2334+
; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r9, %r7, %r5, %r3};
2335+
; CHECK-NEXT: ret;
2336+
%r = uitofp <4 x i8> %a to <4 x float>
2337+
ret <4 x float> %r
2338+
}
2339+
2340+
define <4 x float> @test_sitofp_v4i8(<4 x i8> %a) {
2341+
; CHECK-LABEL: test_sitofp_v4i8(
2342+
; CHECK: {
2343+
; CHECK-NEXT: .reg .b16 %rs<5>;
2344+
; CHECK-NEXT: .reg .b32 %r<10>;
2345+
; CHECK-EMPTY:
2346+
; CHECK-NEXT: // %bb.0:
2347+
; CHECK-NEXT: ld.param.b32 %r1, [test_sitofp_v4i8_param_0];
2348+
; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0xbbb3U;
2349+
; CHECK-NEXT: cvt.u16.u32 %rs1, %r2;
2350+
; CHECK-NEXT: cvt.rn.f32.s16 %r3, %rs1;
2351+
; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0xaaa2U;
2352+
; CHECK-NEXT: cvt.u16.u32 %rs2, %r4;
2353+
; CHECK-NEXT: cvt.rn.f32.s16 %r5, %rs2;
2354+
; CHECK-NEXT: prmt.b32 %r6, %r1, 0, 0x9991U;
2355+
; CHECK-NEXT: cvt.u16.u32 %rs3, %r6;
2356+
; CHECK-NEXT: cvt.rn.f32.s16 %r7, %rs3;
2357+
; CHECK-NEXT: prmt.b32 %r8, %r1, 0, 0x8880U;
2358+
; CHECK-NEXT: cvt.u16.u32 %rs4, %r8;
2359+
; CHECK-NEXT: cvt.rn.f32.s16 %r9, %rs4;
2360+
; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r9, %r7, %r5, %r3};
2361+
; CHECK-NEXT: ret;
2362+
%r = sitofp <4 x i8> %a to <4 x float>
2363+
ret <4 x float> %r
2364+
}
2365+
23142366
attributes #0 = { nounwind }

‎llvm/test/CodeGen/NVPTX/trunc-tofp.ll‎

Lines changed: 85 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,85 @@
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2+
; RUN: llc < %s -mcpu=sm_80 | FileCheck %s
3+
; RUN: %if ptxas %{ llc < %s -mcpu=sm_80 | %ptxas-verify %}
4+
5+
target triple = "nvptx64-nvidia-cuda"
6+
7+
define float @uitofp_trunc_nuw(i32 %x, i32 %y) {
8+
; CHECK-LABEL: uitofp_trunc_nuw(
9+
; CHECK: {
10+
; CHECK-NEXT: .reg .b16 %rs<2>;
11+
; CHECK-NEXT: .reg .b32 %r<5>;
12+
; CHECK-EMPTY:
13+
; CHECK-NEXT: // %bb.0:
14+
; CHECK-NEXT: ld.param.b32 %r1, [uitofp_trunc_nuw_param_0];
15+
; CHECK-NEXT: ld.param.b32 %r2, [uitofp_trunc_nuw_param_1];
16+
; CHECK-NEXT: add.s32 %r3, %r1, %r2;
17+
; CHECK-NEXT: cvt.u16.u32 %rs1, %r3;
18+
; CHECK-NEXT: cvt.rn.f32.u16 %r4, %rs1;
19+
; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
20+
; CHECK-NEXT: ret;
21+
%v = add i32 %x, %y
22+
%t = trunc nuw i32 %v to i16
23+
%f = uitofp i16 %t to float
24+
ret float %f
25+
}
26+
27+
define float @sitofp_trunc_nsw(i32 %x, i32 %y) {
28+
; CHECK-LABEL: sitofp_trunc_nsw(
29+
; CHECK: {
30+
; CHECK-NEXT: .reg .b16 %rs<2>;
31+
; CHECK-NEXT: .reg .b32 %r<5>;
32+
; CHECK-EMPTY:
33+
; CHECK-NEXT: // %bb.0:
34+
; CHECK-NEXT: ld.param.b32 %r1, [sitofp_trunc_nsw_param_0];
35+
; CHECK-NEXT: ld.param.b32 %r2, [sitofp_trunc_nsw_param_1];
36+
; CHECK-NEXT: add.s32 %r3, %r1, %r2;
37+
; CHECK-NEXT: cvt.u16.u32 %rs1, %r3;
38+
; CHECK-NEXT: cvt.rn.f32.s16 %r4, %rs1;
39+
; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
40+
; CHECK-NEXT: ret;
41+
%v = add i32 %x, %y
42+
%t = trunc nsw i32 %v to i16
43+
%f = sitofp i16 %t to float
44+
ret float %f
45+
}
46+
47+
define float @uitofp_trunc_nsw(i32 %x, i32 %y) {
48+
; CHECK-LABEL: uitofp_trunc_nsw(
49+
; CHECK: {
50+
; CHECK-NEXT: .reg .b16 %rs<2>;
51+
; CHECK-NEXT: .reg .b32 %r<5>;
52+
; CHECK-EMPTY:
53+
; CHECK-NEXT: // %bb.0:
54+
; CHECK-NEXT: ld.param.b32 %r1, [uitofp_trunc_nsw_param_0];
55+
; CHECK-NEXT: ld.param.b32 %r2, [uitofp_trunc_nsw_param_1];
56+
; CHECK-NEXT: add.s32 %r3, %r1, %r2;
57+
; CHECK-NEXT: cvt.u16.u32 %rs1, %r3;
58+
; CHECK-NEXT: cvt.rn.f32.u16 %r4, %rs1;
59+
; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
60+
; CHECK-NEXT: ret;
61+
%v = add i32 %x, %y
62+
%t = trunc nsw i32 %v to i16
63+
%f = uitofp i16 %t to float
64+
ret float %f
65+
}
66+
67+
define float @sitofp_trunc_nuw(i32 %x, i32 %y) {
68+
; CHECK-LABEL: sitofp_trunc_nuw(
69+
; CHECK: {
70+
; CHECK-NEXT: .reg .b16 %rs<2>;
71+
; CHECK-NEXT: .reg .b32 %r<5>;
72+
; CHECK-EMPTY:
73+
; CHECK-NEXT: // %bb.0:
74+
; CHECK-NEXT: ld.param.b32 %r1, [sitofp_trunc_nuw_param_0];
75+
; CHECK-NEXT: ld.param.b32 %r2, [sitofp_trunc_nuw_param_1];
76+
; CHECK-NEXT: add.s32 %r3, %r1, %r2;
77+
; CHECK-NEXT: cvt.u16.u32 %rs1, %r3;
78+
; CHECK-NEXT: cvt.rn.f32.s16 %r4, %rs1;
79+
; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
80+
; CHECK-NEXT: ret;
81+
%v = add i32 %x, %y
82+
%t = trunc nuw i32 %v to i16
83+
%f = sitofp i16 %t to float
84+
ret float %f
85+
}

0 commit comments

Comments
 (0)