diff options
author | Alex MacLean <amaclean@nvidia.com> | 2024-01-17 16:18:39 -0800 |
---|---|---|
committer | GitHub <noreply@github.com> | 2024-01-17 16:18:39 -0800 |
commit | 430a40d12eaa5a61792c4670955c110146902afb (patch) | |
tree | 33f82c05a67529ee9aa3efbad1d7d3b8aa610ba4 /llvm/test/CodeGen | |
parent | f2b5a314b29275f2092af3ec26f42272daa4312c (diff) | |
download | llvm-430a40d12eaa5a61792c4670955c110146902afb.zip llvm-430a40d12eaa5a61792c4670955c110146902afb.tar.gz llvm-430a40d12eaa5a61792c4670955c110146902afb.tar.bz2 |
[NVPTX] extend type support for nvvm.{min,max,mulhi,sad} (#78385)
Ensure intrinsics and auto-upgrades support i16, i32, and i64 for for
`nvvm.{min,max,mulhi,sad}`
- `nvvm.min` and `nvvm.max`: These are auto-upgraded to `select`
instructions but it is still nice to support the 16 bit variants just in
case any generators of IR are still trying to use these intrinsics.
- `nvvm.sad` added both the 16 and 64 bit variants, also marked this
instruction as speculateble. These directly correspond to the PTX
`sad.{u16,s16,u64,s64}` instructions.
- `nvvm.mulhi` added the 16 bit variants. These directly correspond to
the PTX `mul.hi.{s,u}16` instructions.
Diffstat (limited to 'llvm/test/CodeGen')
-rw-r--r-- | llvm/test/CodeGen/NVPTX/mulhi-intrins.ll | 104 | ||||
-rw-r--r-- | llvm/test/CodeGen/NVPTX/sad-intrins.ll | 110 |
2 files changed, 214 insertions, 0 deletions
diff --git a/llvm/test/CodeGen/NVPTX/mulhi-intrins.ll b/llvm/test/CodeGen/NVPTX/mulhi-intrins.ll new file mode 100644 index 0000000..efa9946 --- /dev/null +++ b/llvm/test/CodeGen/NVPTX/mulhi-intrins.ll @@ -0,0 +1,104 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 4 +; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_50 | FileCheck %s +; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_50 | %ptxas-verify %} + +define i16 @test_mulhi_i16(i16 %x, i16 %y) { +; CHECK-LABEL: test_mulhi_i16( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<4>; +; CHECK-NEXT: .reg .b32 %r<2>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.u16 %rs1, [test_mulhi_i16_param_0]; +; CHECK-NEXT: ld.param.u16 %rs2, [test_mulhi_i16_param_1]; +; CHECK-NEXT: mul.hi.s16 %rs3, %rs1, %rs2; +; CHECK-NEXT: cvt.u32.u16 %r1, %rs3; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r1; +; CHECK-NEXT: ret; + %1 = call i16 @llvm.nvvm.mulhi.s(i16 %x, i16 %y) + ret i16 %1 +} + +define i16 @test_mulhi_u16(i16 %x, i16 %y) { +; CHECK-LABEL: test_mulhi_u16( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<4>; +; CHECK-NEXT: .reg .b32 %r<2>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.u16 %rs1, [test_mulhi_u16_param_0]; +; CHECK-NEXT: ld.param.u16 %rs2, [test_mulhi_u16_param_1]; +; CHECK-NEXT: mul.hi.u16 %rs3, %rs1, %rs2; +; CHECK-NEXT: cvt.u32.u16 %r1, %rs3; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r1; +; CHECK-NEXT: ret; + %1 = call i16 @llvm.nvvm.mulhi.us(i16 %x, i16 %y) + ret i16 %1 +} + +define i32 @test_mulhi_i32(i32 %x, i32 %y) { +; CHECK-LABEL: test_mulhi_i32( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.u32 %r1, [test_mulhi_i32_param_0]; +; CHECK-NEXT: ld.param.u32 %r2, [test_mulhi_i32_param_1]; +; CHECK-NEXT: mul.hi.s32 %r3, %r1, %r2; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3; +; CHECK-NEXT: ret; + %1 = call i32 @llvm.nvvm.mulhi.i(i32 %x, i32 %y) + ret i32 %1 +} + +define i32 @test_mulhi_u32(i32 %x, i32 %y) { +; CHECK-LABEL: test_mulhi_u32( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.u32 %r1, [test_mulhi_u32_param_0]; +; CHECK-NEXT: ld.param.u32 %r2, [test_mulhi_u32_param_1]; +; CHECK-NEXT: mul.hi.u32 %r3, %r1, %r2; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3; +; CHECK-NEXT: ret; + %1 = call i32 @llvm.nvvm.mulhi.ui(i32 %x, i32 %y) + ret i32 %1 +} + +define i64 @test_mulhi_i64(i64 %x, i64 %y) { +; CHECK-LABEL: test_mulhi_i64( +; CHECK: { +; CHECK-NEXT: .reg .b64 %rd<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.u64 %rd1, [test_mulhi_i64_param_0]; +; CHECK-NEXT: ld.param.u64 %rd2, [test_mulhi_i64_param_1]; +; CHECK-NEXT: mul.hi.s64 %rd3, %rd1, %rd2; +; CHECK-NEXT: st.param.b64 [func_retval0+0], %rd3; +; CHECK-NEXT: ret; + %1 = call i64 @llvm.nvvm.mulhi.ll(i64 %x, i64 %y) + ret i64 %1 +} + +define i64 @test_mulhi_u64(i64 %x, i64 %y) { +; CHECK-LABEL: test_mulhi_u64( +; CHECK: { +; CHECK-NEXT: .reg .b64 %rd<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.u64 %rd1, [test_mulhi_u64_param_0]; +; CHECK-NEXT: ld.param.u64 %rd2, [test_mulhi_u64_param_1]; +; CHECK-NEXT: mul.hi.u64 %rd3, %rd1, %rd2; +; CHECK-NEXT: st.param.b64 [func_retval0+0], %rd3; +; CHECK-NEXT: ret; + %1 = call i64 @llvm.nvvm.mulhi.ull(i64 %x, i64 %y) + ret i64 %1 +} + +declare i16 @llvm.nvvm.mulhi.s(i16, i16) +declare i16 @llvm.nvvm.mulhi.us(i16, i16) +declare i32 @llvm.nvvm.mulhi.i(i32, i32) +declare i32 @llvm.nvvm.mulhi.ui(i32, i32) +declare i64 @llvm.nvvm.mulhi.ll(i64, i64) +declare i64 @llvm.nvvm.mulhi.ull(i64, i64) diff --git a/llvm/test/CodeGen/NVPTX/sad-intrins.ll b/llvm/test/CodeGen/NVPTX/sad-intrins.ll new file mode 100644 index 0000000..a09413b --- /dev/null +++ b/llvm/test/CodeGen/NVPTX/sad-intrins.ll @@ -0,0 +1,110 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 4 +; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_50 | FileCheck %s +; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_50 | %ptxas-verify %} + +define i16 @test_sad_i16(i16 %x, i16 %y, i16 %z) { +; CHECK-LABEL: test_sad_i16( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<5>; +; CHECK-NEXT: .reg .b32 %r<2>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.u16 %rs1, [test_sad_i16_param_0]; +; CHECK-NEXT: ld.param.u16 %rs2, [test_sad_i16_param_1]; +; CHECK-NEXT: ld.param.u16 %rs3, [test_sad_i16_param_2]; +; CHECK-NEXT: sad.s16 %rs4, %rs1, %rs2, %rs3; +; CHECK-NEXT: cvt.u32.u16 %r1, %rs4; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r1; +; CHECK-NEXT: ret; + %1 = call i16 @llvm.nvvm.sad.s(i16 %x, i16 %y, i16 %z) + ret i16 %1 +} + +define i16 @test_sad_u16(i16 %x, i16 %y, i16 %z) { +; CHECK-LABEL: test_sad_u16( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<5>; +; CHECK-NEXT: .reg .b32 %r<2>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.u16 %rs1, [test_sad_u16_param_0]; +; CHECK-NEXT: ld.param.u16 %rs2, [test_sad_u16_param_1]; +; CHECK-NEXT: ld.param.u16 %rs3, [test_sad_u16_param_2]; +; CHECK-NEXT: sad.u16 %rs4, %rs1, %rs2, %rs3; +; CHECK-NEXT: cvt.u32.u16 %r1, %rs4; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r1; +; CHECK-NEXT: ret; + %1 = call i16 @llvm.nvvm.sad.us(i16 %x, i16 %y, i16 %z) + ret i16 %1 +} + +define i32 @test_sad_i32(i32 %x, i32 %y, i32 %z) { +; CHECK-LABEL: test_sad_i32( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<5>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.u32 %r1, [test_sad_i32_param_0]; +; CHECK-NEXT: ld.param.u32 %r2, [test_sad_i32_param_1]; +; CHECK-NEXT: ld.param.u32 %r3, [test_sad_i32_param_2]; +; CHECK-NEXT: sad.s32 %r4, %r1, %r2, %r3; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r4; +; CHECK-NEXT: ret; + %1 = call i32 @llvm.nvvm.sad.i(i32 %x, i32 %y, i32 %z) + ret i32 %1 +} + +define i32 @test_sad_u32(i32 %x, i32 %y, i32 %z) { +; CHECK-LABEL: test_sad_u32( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<5>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.u32 %r1, [test_sad_u32_param_0]; +; CHECK-NEXT: ld.param.u32 %r2, [test_sad_u32_param_1]; +; CHECK-NEXT: ld.param.u32 %r3, [test_sad_u32_param_2]; +; CHECK-NEXT: sad.u32 %r4, %r1, %r2, %r3; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r4; +; CHECK-NEXT: ret; + %1 = call i32 @llvm.nvvm.sad.ui(i32 %x, i32 %y, i32 %z) + ret i32 %1 +} + +define i64 @test_sad_i64(i64 %x, i64 %y, i64 %z) { +; CHECK-LABEL: test_sad_i64( +; CHECK: { +; CHECK-NEXT: .reg .b64 %rd<5>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.u64 %rd1, [test_sad_i64_param_0]; +; CHECK-NEXT: ld.param.u64 %rd2, [test_sad_i64_param_1]; +; CHECK-NEXT: ld.param.u64 %rd3, [test_sad_i64_param_2]; +; CHECK-NEXT: sad.s64 %rd4, %rd1, %rd2, %rd3; +; CHECK-NEXT: st.param.b64 [func_retval0+0], %rd4; +; CHECK-NEXT: ret; + %1 = call i64 @llvm.nvvm.sad.ll(i64 %x, i64 %y, i64 %z) + ret i64 %1 +} + +define i64 @test_sad_u64(i64 %x, i64 %y, i64 %z) { +; CHECK-LABEL: test_sad_u64( +; CHECK: { +; CHECK-NEXT: .reg .b64 %rd<5>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.u64 %rd1, [test_sad_u64_param_0]; +; CHECK-NEXT: ld.param.u64 %rd2, [test_sad_u64_param_1]; +; CHECK-NEXT: ld.param.u64 %rd3, [test_sad_u64_param_2]; +; CHECK-NEXT: sad.u64 %rd4, %rd1, %rd2, %rd3; +; CHECK-NEXT: st.param.b64 [func_retval0+0], %rd4; +; CHECK-NEXT: ret; + %1 = call i64 @llvm.nvvm.sad.ull(i64 %x, i64 %y, i64 %z) + ret i64 %1 +} + +declare i16 @llvm.nvvm.sad.s(i16, i16, i16) +declare i16 @llvm.nvvm.sad.us(i16, i16, i16) +declare i32 @llvm.nvvm.sad.i(i32, i32, i32) +declare i32 @llvm.nvvm.sad.ui(i32, i32, i32) +declare i64 @llvm.nvvm.sad.ll(i64, i64, i64) +declare i64 @llvm.nvvm.sad.ull(i64, i64, i64) |