131 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -verify-machineinstrs -mcpu=sm_50 | FileCheck %s3; RUN: %if ptxas %{ llc < %s -verify-machineinstrs -mcpu=sm_50 | %ptxas-verify %}4 5target triple = "nvptx64-nvidia-cuda"6 7define i32 @test_prmt_basic(i32 %lo, i32 %hi, i32 %selector) {8; CHECK-LABEL: test_prmt_basic(9; CHECK: {10; CHECK-NEXT: .reg .b32 %r<5>;11; CHECK-EMPTY:12; CHECK-NEXT: // %bb.0:13; CHECK-NEXT: ld.param.b32 %r1, [test_prmt_basic_param_0];14; CHECK-NEXT: ld.param.b32 %r2, [test_prmt_basic_param_1];15; CHECK-NEXT: ld.param.b32 %r3, [test_prmt_basic_param_2];16; CHECK-NEXT: prmt.b32 %r4, %r1, %r2, %r3;17; CHECK-NEXT: st.param.b32 [func_retval0], %r4;18; CHECK-NEXT: ret;19 %val = call i32 @llvm.nvvm.prmt(i32 %lo, i32 %hi, i32 %selector)20 ret i32 %val21}22 23define i32 @test_prmt_f4e(i32 %lo, i32 %hi, i32 %selector) {24; CHECK-LABEL: test_prmt_f4e(25; CHECK: {26; CHECK-NEXT: .reg .b32 %r<5>;27; CHECK-EMPTY:28; CHECK-NEXT: // %bb.0:29; CHECK-NEXT: ld.param.b32 %r1, [test_prmt_f4e_param_0];30; CHECK-NEXT: ld.param.b32 %r2, [test_prmt_f4e_param_1];31; CHECK-NEXT: ld.param.b32 %r3, [test_prmt_f4e_param_2];32; CHECK-NEXT: prmt.b32.f4e %r4, %r1, %r2, %r3;33; CHECK-NEXT: st.param.b32 [func_retval0], %r4;34; CHECK-NEXT: ret;35 %val = call i32 @llvm.nvvm.prmt.f4e(i32 %lo, i32 %hi, i32 %selector)36 ret i32 %val37}38 39define i32 @test_prmt_b4e(i32 %lo, i32 %hi, i32 %selector) {40; CHECK-LABEL: test_prmt_b4e(41; CHECK: {42; CHECK-NEXT: .reg .b32 %r<5>;43; CHECK-EMPTY:44; CHECK-NEXT: // %bb.0:45; CHECK-NEXT: ld.param.b32 %r1, [test_prmt_b4e_param_0];46; CHECK-NEXT: ld.param.b32 %r2, [test_prmt_b4e_param_1];47; CHECK-NEXT: ld.param.b32 %r3, [test_prmt_b4e_param_2];48; CHECK-NEXT: prmt.b32.b4e %r4, %r1, %r2, %r3;49; CHECK-NEXT: st.param.b32 [func_retval0], %r4;50; CHECK-NEXT: ret;51 %val = call i32 @llvm.nvvm.prmt.b4e(i32 %lo, i32 %hi, i32 %selector)52 ret i32 %val53}54 55define i32 @test_prmt_rc8(i32 %lo, i32 %selector) {56; CHECK-LABEL: test_prmt_rc8(57; CHECK: {58; CHECK-NEXT: .reg .b32 %r<4>;59; CHECK-EMPTY:60; CHECK-NEXT: // %bb.0:61; CHECK-NEXT: ld.param.b32 %r1, [test_prmt_rc8_param_0];62; CHECK-NEXT: ld.param.b32 %r2, [test_prmt_rc8_param_1];63; CHECK-NEXT: prmt.b32.rc8 %r3, %r1, 0, %r2;64; CHECK-NEXT: st.param.b32 [func_retval0], %r3;65; CHECK-NEXT: ret;66 %val = call i32 @llvm.nvvm.prmt.rc8(i32 %lo, i32 %selector)67 ret i32 %val68}69 70define i32 @test_prmt_ecl(i32 %lo, i32 %selector) {71; CHECK-LABEL: test_prmt_ecl(72; CHECK: {73; CHECK-NEXT: .reg .b32 %r<4>;74; CHECK-EMPTY:75; CHECK-NEXT: // %bb.0:76; CHECK-NEXT: ld.param.b32 %r1, [test_prmt_ecl_param_0];77; CHECK-NEXT: ld.param.b32 %r2, [test_prmt_ecl_param_1];78; CHECK-NEXT: prmt.b32.ecl %r3, %r1, 0, %r2;79; CHECK-NEXT: st.param.b32 [func_retval0], %r3;80; CHECK-NEXT: ret;81 %val = call i32 @llvm.nvvm.prmt.ecl(i32 %lo, i32 %selector)82 ret i32 %val83}84 85define i32 @test_prmt_ecr(i32 %lo, i32 %selector) {86; CHECK-LABEL: test_prmt_ecr(87; CHECK: {88; CHECK-NEXT: .reg .b32 %r<4>;89; CHECK-EMPTY:90; CHECK-NEXT: // %bb.0:91; CHECK-NEXT: ld.param.b32 %r1, [test_prmt_ecr_param_0];92; CHECK-NEXT: ld.param.b32 %r2, [test_prmt_ecr_param_1];93; CHECK-NEXT: prmt.b32.ecr %r3, %r1, 0, %r2;94; CHECK-NEXT: st.param.b32 [func_retval0], %r3;95; CHECK-NEXT: ret;96 %val = call i32 @llvm.nvvm.prmt.ecr(i32 %lo, i32 %selector)97 ret i32 %val98}99 100define i32 @test_prmt_rc16(i32 %lo, i32 %selector) {101; CHECK-LABEL: test_prmt_rc16(102; CHECK: {103; CHECK-NEXT: .reg .b32 %r<4>;104; CHECK-EMPTY:105; CHECK-NEXT: // %bb.0:106; CHECK-NEXT: ld.param.b32 %r1, [test_prmt_rc16_param_0];107; CHECK-NEXT: ld.param.b32 %r2, [test_prmt_rc16_param_1];108; CHECK-NEXT: prmt.b32.rc16 %r3, %r1, 0, %r2;109; CHECK-NEXT: st.param.b32 [func_retval0], %r3;110; CHECK-NEXT: ret;111 %val = call i32 @llvm.nvvm.prmt.rc16(i32 %lo, i32 %selector)112 ret i32 %val113}114 115define i32 @test_prmt_f4e_folding(i32 %lo, i32 %hi, i32 %ptr) {116; CHECK-LABEL: test_prmt_f4e_folding(117; CHECK: {118; CHECK-NEXT: .reg .b32 %r<5>;119; CHECK-EMPTY:120; CHECK-NEXT: // %bb.0:121; CHECK-NEXT: ld.param.b32 %r1, [test_prmt_f4e_folding_param_0];122; CHECK-NEXT: ld.param.b32 %r2, [test_prmt_f4e_folding_param_1];123; CHECK-NEXT: ld.param.b32 %r3, [test_prmt_f4e_folding_param_2];124; CHECK-NEXT: prmt.b32.f4e %r4, %r1, %r2, %r3;125; CHECK-NEXT: st.param.b32 [func_retval0], %r4;126; CHECK-NEXT: ret;127 %sh_amt = shl i32 %ptr, 3128 %val = call i32 @llvm.fshr.i32(i32 %hi, i32 %lo, i32 %sh_amt)129 ret i32 %val130}131