338 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx -mcpu=sm_60 | FileCheck %s --check-prefixes=CHECK,CHECK323; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_60 | FileCheck %s --check-prefixes=CHECK,CHECK644; RUN: %if ptxas-sm_60 && ptxas-ptr32 %{ llc < %s -mtriple=nvptx -mcpu=sm_60 | %ptxas-verify -arch=sm_60 %}5; RUN: %if ptxas-sm_60 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_60 | %ptxas-verify -arch=sm_60 %}6 7define float @test_fabsf(float %f) {8; CHECK-LABEL: test_fabsf(9; CHECK: {10; CHECK-NEXT: .reg .b32 %r<3>;11; CHECK-EMPTY:12; CHECK-NEXT: // %bb.0:13; CHECK-NEXT: ld.param.b32 %r1, [test_fabsf_param_0];14; CHECK-NEXT: abs.f32 %r2, %r1;15; CHECK-NEXT: st.param.b32 [func_retval0], %r2;16; CHECK-NEXT: ret;17 %x = call float @llvm.fabs.f32(float %f)18 ret float %x19}20 21define double @test_fabs(double %d) {22; CHECK-LABEL: test_fabs(23; CHECK: {24; CHECK-NEXT: .reg .b64 %rd<3>;25; CHECK-EMPTY:26; CHECK-NEXT: // %bb.0:27; CHECK-NEXT: ld.param.b64 %rd1, [test_fabs_param_0];28; CHECK-NEXT: abs.f64 %rd2, %rd1;29; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;30; CHECK-NEXT: ret;31 %x = call double @llvm.fabs.f64(double %d)32 ret double %x33}34 35define float @test_nvvm_sqrt(float %a) {36; CHECK-LABEL: test_nvvm_sqrt(37; CHECK: {38; CHECK-NEXT: .reg .b32 %r<3>;39; CHECK-EMPTY:40; CHECK-NEXT: // %bb.0:41; CHECK-NEXT: ld.param.b32 %r1, [test_nvvm_sqrt_param_0];42; CHECK-NEXT: sqrt.rn.f32 %r2, %r1;43; CHECK-NEXT: st.param.b32 [func_retval0], %r2;44; CHECK-NEXT: ret;45 %val = call float @llvm.nvvm.sqrt.f(float %a)46 ret float %val47}48 49define float @test_llvm_sqrt(float %a) {50; CHECK-LABEL: test_llvm_sqrt(51; CHECK: {52; CHECK-NEXT: .reg .b32 %r<3>;53; CHECK-EMPTY:54; CHECK-NEXT: // %bb.0:55; CHECK-NEXT: ld.param.b32 %r1, [test_llvm_sqrt_param_0];56; CHECK-NEXT: sqrt.rn.f32 %r2, %r1;57; CHECK-NEXT: st.param.b32 [func_retval0], %r2;58; CHECK-NEXT: ret;59 %val = call float @llvm.sqrt.f32(float %a)60 ret float %val61}62 63define i32 @test_bitreverse32(i32 %a) {64; CHECK-LABEL: test_bitreverse32(65; CHECK: {66; CHECK-NEXT: .reg .b32 %r<3>;67; CHECK-EMPTY:68; CHECK-NEXT: // %bb.0:69; CHECK-NEXT: ld.param.b32 %r1, [test_bitreverse32_param_0];70; CHECK-NEXT: brev.b32 %r2, %r1;71; CHECK-NEXT: st.param.b32 [func_retval0], %r2;72; CHECK-NEXT: ret;73 %val = call i32 @llvm.bitreverse.i32(i32 %a)74 ret i32 %val75}76 77define i64 @test_bitreverse64(i64 %a) {78; CHECK-LABEL: test_bitreverse64(79; CHECK: {80; CHECK-NEXT: .reg .b64 %rd<3>;81; CHECK-EMPTY:82; CHECK-NEXT: // %bb.0:83; CHECK-NEXT: ld.param.b64 %rd1, [test_bitreverse64_param_0];84; CHECK-NEXT: brev.b64 %rd2, %rd1;85; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;86; CHECK-NEXT: ret;87 %val = call i64 @llvm.bitreverse.i64(i64 %a)88 ret i64 %val89}90 91define i32 @test_popc32(i32 %a) {92; CHECK-LABEL: test_popc32(93; CHECK: {94; CHECK-NEXT: .reg .b32 %r<3>;95; CHECK-EMPTY:96; CHECK-NEXT: // %bb.0:97; CHECK-NEXT: ld.param.b32 %r1, [test_popc32_param_0];98; CHECK-NEXT: popc.b32 %r2, %r1;99; CHECK-NEXT: st.param.b32 [func_retval0], %r2;100; CHECK-NEXT: ret;101 %val = call i32 @llvm.ctpop.i32(i32 %a)102 ret i32 %val103}104 105define i64 @test_popc64(i64 %a) {106; CHECK-LABEL: test_popc64(107; CHECK: {108; CHECK-NEXT: .reg .b32 %r<2>;109; CHECK-NEXT: .reg .b64 %rd<3>;110; CHECK-EMPTY:111; CHECK-NEXT: // %bb.0:112; CHECK-NEXT: ld.param.b64 %rd1, [test_popc64_param_0];113; CHECK-NEXT: popc.b64 %r1, %rd1;114; CHECK-NEXT: cvt.u64.u32 %rd2, %r1;115; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;116; CHECK-NEXT: ret;117 %val = call i64 @llvm.ctpop.i64(i64 %a)118 ret i64 %val119}120 121; NVPTX popc.b64 returns an i32 even though @llvm.ctpop.i64 returns an i64, so122; if this function returns an i32, there's no need to do any type conversions123; in the ptx.124define i32 @test_popc64_trunc(i64 %a) {125; CHECK-LABEL: test_popc64_trunc(126; CHECK: {127; CHECK-NEXT: .reg .b32 %r<2>;128; CHECK-NEXT: .reg .b64 %rd<2>;129; CHECK-EMPTY:130; CHECK-NEXT: // %bb.0:131; CHECK-NEXT: ld.param.b64 %rd1, [test_popc64_trunc_param_0];132; CHECK-NEXT: popc.b64 %r1, %rd1;133; CHECK-NEXT: st.param.b32 [func_retval0], %r1;134; CHECK-NEXT: ret;135 %val = call i64 @llvm.ctpop.i64(i64 %a)136 %trunc = trunc i64 %val to i32137 ret i32 %trunc138}139 140; llvm.ctpop.i16 is implemenented by converting to i32, running popc.b32, and141; then converting back to i16.142define void @test_popc16(i16 %a, ptr %b) {143; CHECK32-LABEL: test_popc16(144; CHECK32: {145; CHECK32-NEXT: .reg .b32 %r<4>;146; CHECK32-EMPTY:147; CHECK32-NEXT: // %bb.0:148; CHECK32-NEXT: ld.param.b16 %r1, [test_popc16_param_0];149; CHECK32-NEXT: popc.b32 %r2, %r1;150; CHECK32-NEXT: ld.param.b32 %r3, [test_popc16_param_1];151; CHECK32-NEXT: st.b16 [%r3], %r2;152; CHECK32-NEXT: ret;153;154; CHECK64-LABEL: test_popc16(155; CHECK64: {156; CHECK64-NEXT: .reg .b32 %r<3>;157; CHECK64-NEXT: .reg .b64 %rd<2>;158; CHECK64-EMPTY:159; CHECK64-NEXT: // %bb.0:160; CHECK64-NEXT: ld.param.b16 %r1, [test_popc16_param_0];161; CHECK64-NEXT: popc.b32 %r2, %r1;162; CHECK64-NEXT: ld.param.b64 %rd1, [test_popc16_param_1];163; CHECK64-NEXT: st.b16 [%rd1], %r2;164; CHECK64-NEXT: ret;165 %val = call i16 @llvm.ctpop.i16(i16 %a)166 store i16 %val, ptr %b167 ret void168}169 170; If we call llvm.ctpop.i16 and then zext the result to i32, we shouldn't need171; to do any conversions after calling popc.b32, because that returns an i32.172define i32 @test_popc16_to_32(i16 %a) {173; CHECK-LABEL: test_popc16_to_32(174; CHECK: {175; CHECK-NEXT: .reg .b32 %r<3>;176; CHECK-EMPTY:177; CHECK-NEXT: // %bb.0:178; CHECK-NEXT: ld.param.b16 %r1, [test_popc16_to_32_param_0];179; CHECK-NEXT: popc.b32 %r2, %r1;180; CHECK-NEXT: st.param.b32 [func_retval0], %r2;181; CHECK-NEXT: ret;182 %val = call i16 @llvm.ctpop.i16(i16 %a)183 %zext = zext i16 %val to i32184 ret i32 %zext185}186 187; Most of nvvm.read.ptx.sreg.* intrinsics always return the same value and may188; be CSE'd.189define i32 @test_tid() {190; CHECK-LABEL: test_tid(191; CHECK: {192; CHECK-NEXT: .reg .b32 %r<3>;193; CHECK-EMPTY:194; CHECK-NEXT: // %bb.0:195; CHECK-NEXT: mov.u32 %r1, %tid.x;196; CHECK-NEXT: add.s32 %r2, %r1, %r1;197; CHECK-NEXT: st.param.b32 [func_retval0], %r2;198; CHECK-NEXT: ret;199 %a = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x()200 %b = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x()201 %ret = add i32 %a, %b202 ret i32 %ret203}204 205; reading clock() or clock64() should not be CSE'd as each read may return206; different value.207define i32 @test_clock() {208; CHECK-LABEL: test_clock(209; CHECK: {210; CHECK-NEXT: .reg .b32 %r<4>;211; CHECK-EMPTY:212; CHECK-NEXT: // %bb.0:213; CHECK-NEXT: mov.u32 %r1, %clock;214; CHECK-NEXT: mov.u32 %r2, %clock;215; CHECK-NEXT: add.s32 %r3, %r1, %r2;216; CHECK-NEXT: st.param.b32 [func_retval0], %r3;217; CHECK-NEXT: ret;218 %a = tail call i32 @llvm.nvvm.read.ptx.sreg.clock()219 %b = tail call i32 @llvm.nvvm.read.ptx.sreg.clock()220 %ret = add i32 %a, %b221 ret i32 %ret222}223 224define i64 @test_clock64() {225; CHECK-LABEL: test_clock64(226; CHECK: {227; CHECK-NEXT: .reg .b64 %rd<4>;228; CHECK-EMPTY:229; CHECK-NEXT: // %bb.0:230; CHECK-NEXT: mov.u64 %rd1, %clock64;231; CHECK-NEXT: mov.u64 %rd2, %clock64;232; CHECK-NEXT: add.s64 %rd3, %rd1, %rd2;233; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;234; CHECK-NEXT: ret;235 %a = tail call i64 @llvm.nvvm.read.ptx.sreg.clock64()236 %b = tail call i64 @llvm.nvvm.read.ptx.sreg.clock64()237 %ret = add i64 %a, %b238 ret i64 %ret239}240 241define void @test_exit() {242; CHECK-LABEL: test_exit(243; CHECK: {244; CHECK-EMPTY:245; CHECK-EMPTY:246; CHECK-NEXT: // %bb.0:247; CHECK-NEXT: exit;248; CHECK-NEXT: ret;249 call void @llvm.nvvm.exit()250 ret void251}252 253define i64 @test_globaltimer() {254; CHECK-LABEL: test_globaltimer(255; CHECK: {256; CHECK-NEXT: .reg .b64 %rd<4>;257; CHECK-EMPTY:258; CHECK-NEXT: // %bb.0:259; CHECK-NEXT: mov.u64 %rd1, %globaltimer;260; CHECK-NEXT: mov.u64 %rd2, %globaltimer;261; CHECK-NEXT: add.s64 %rd3, %rd1, %rd2;262; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;263; CHECK-NEXT: ret;264 %a = tail call i64 @llvm.nvvm.read.ptx.sreg.globaltimer()265 %b = tail call i64 @llvm.nvvm.read.ptx.sreg.globaltimer()266 %ret = add i64 %a, %b267 ret i64 %ret268}269 270define i32 @test_globaltimer_lo(){271; CHECK-LABEL: test_globaltimer_lo(272; CHECK: {273; CHECK-NEXT: .reg .b32 %r<4>;274; CHECK-EMPTY:275; CHECK-NEXT: // %bb.0:276; CHECK-NEXT: mov.u32 %r1, %globaltimer_lo;277; CHECK-NEXT: mov.u32 %r2, %globaltimer_lo;278; CHECK-NEXT: add.s32 %r3, %r1, %r2;279; CHECK-NEXT: st.param.b32 [func_retval0], %r3;280; CHECK-NEXT: ret;281 %a = tail call i32 @llvm.nvvm.read.ptx.sreg.globaltimer.lo()282 %b = tail call i32 @llvm.nvvm.read.ptx.sreg.globaltimer.lo()283 %ret = add i32 %a, %b284 ret i32 %ret285}286 287define i64 @test_cyclecounter() {288; CHECK-LABEL: test_cyclecounter(289; CHECK: {290; CHECK-NEXT: .reg .b64 %rd<4>;291; CHECK-EMPTY:292; CHECK-NEXT: // %bb.0:293; CHECK-NEXT: mov.u64 %rd1, %clock64;294; CHECK-NEXT: mov.u64 %rd2, %clock64;295; CHECK-NEXT: add.s64 %rd3, %rd1, %rd2;296; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;297; CHECK-NEXT: ret;298 %a = tail call i64 @llvm.readcyclecounter()299 %b = tail call i64 @llvm.readcyclecounter()300 %ret = add i64 %a, %b301 ret i64 %ret302}303 304define i64 @test_steadycounter() {305; CHECK-LABEL: test_steadycounter(306; CHECK: {307; CHECK-NEXT: .reg .b64 %rd<4>;308; CHECK-EMPTY:309; CHECK-NEXT: // %bb.0:310; CHECK-NEXT: mov.u64 %rd1, %globaltimer;311; CHECK-NEXT: mov.u64 %rd2, %globaltimer;312; CHECK-NEXT: add.s64 %rd3, %rd1, %rd2;313; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;314; CHECK-NEXT: ret;315 %a = tail call i64 @llvm.readsteadycounter()316 %b = tail call i64 @llvm.readsteadycounter()317 %ret = add i64 %a, %b318 ret i64 %ret319}320 321declare float @llvm.fabs.f32(float)322declare double @llvm.fabs.f64(double)323declare float @llvm.nvvm.sqrt.f(float)324declare float @llvm.sqrt.f32(float)325declare i32 @llvm.bitreverse.i32(i32)326declare i64 @llvm.bitreverse.i64(i64)327declare i16 @llvm.ctpop.i16(i16)328declare i32 @llvm.ctpop.i32(i32)329declare i64 @llvm.ctpop.i64(i64)330 331declare i32 @llvm.nvvm.read.ptx.sreg.tid.x()332declare i32 @llvm.nvvm.read.ptx.sreg.clock()333declare i64 @llvm.nvvm.read.ptx.sreg.clock64()334declare void @llvm.nvvm.exit()335declare i64 @llvm.nvvm.read.ptx.sreg.globaltimer()336declare i64 @llvm.readcyclecounter()337declare i64 @llvm.readsteadycounter()338