177 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}4 5target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"6 7declare i16 @llvm.ctlz.i16(i16, i1) readnone8declare i32 @llvm.ctlz.i32(i32, i1) readnone9declare i64 @llvm.ctlz.i64(i64, i1) readnone10 11; There should be no difference between llvm.ctlz.i32(%a, true) and12; llvm.ctlz.i32(%a, false), as ptx's clz(0) is defined to return 0.13 14define i32 @myctlz(i32 %a) {15; CHECK-LABEL: myctlz(16; CHECK: {17; CHECK-NEXT: .reg .b32 %r<3>;18; CHECK-EMPTY:19; CHECK-NEXT: // %bb.0:20; CHECK-NEXT: ld.param.b32 %r1, [myctlz_param_0];21; CHECK-NEXT: clz.b32 %r2, %r1;22; CHECK-NEXT: st.param.b32 [func_retval0], %r2;23; CHECK-NEXT: ret;24 %val = call i32 @llvm.ctlz.i32(i32 %a, i1 false) readnone25 ret i32 %val26}27define i32 @myctlz_2(i32 %a) {28; CHECK-LABEL: myctlz_2(29; CHECK: {30; CHECK-NEXT: .reg .b32 %r<3>;31; CHECK-EMPTY:32; CHECK-NEXT: // %bb.0:33; CHECK-NEXT: ld.param.b32 %r1, [myctlz_2_param_0];34; CHECK-NEXT: clz.b32 %r2, %r1;35; CHECK-NEXT: st.param.b32 [func_retval0], %r2;36; CHECK-NEXT: ret;37 %val = call i32 @llvm.ctlz.i32(i32 %a, i1 true) readnone38 ret i32 %val39}40 41; PTX's clz.b64 returns a 32-bit value, but LLVM's intrinsic returns a 64-bit42; value, so here we have to zero-extend it.43define i64 @myctlz64(i64 %a) {44; CHECK-LABEL: myctlz64(45; CHECK: {46; CHECK-NEXT: .reg .b32 %r<2>;47; CHECK-NEXT: .reg .b64 %rd<3>;48; CHECK-EMPTY:49; CHECK-NEXT: // %bb.0:50; CHECK-NEXT: ld.param.b64 %rd1, [myctlz64_param_0];51; CHECK-NEXT: clz.b64 %r1, %rd1;52; CHECK-NEXT: cvt.u64.u32 %rd2, %r1;53; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;54; CHECK-NEXT: ret;55 %val = call i64 @llvm.ctlz.i64(i64 %a, i1 false) readnone56 ret i64 %val57}58define i64 @myctlz64_2(i64 %a) {59; CHECK-LABEL: myctlz64_2(60; CHECK: {61; CHECK-NEXT: .reg .b32 %r<2>;62; CHECK-NEXT: .reg .b64 %rd<3>;63; CHECK-EMPTY:64; CHECK-NEXT: // %bb.0:65; CHECK-NEXT: ld.param.b64 %rd1, [myctlz64_2_param_0];66; CHECK-NEXT: clz.b64 %r1, %rd1;67; CHECK-NEXT: cvt.u64.u32 %rd2, %r1;68; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;69; CHECK-NEXT: ret;70 %val = call i64 @llvm.ctlz.i64(i64 %a, i1 true) readnone71 ret i64 %val72}73 74; Here we truncate the 64-bit value of LLVM's ctlz intrinsic to 32 bits, the75; natural return width of ptx's clz.b64 instruction. No conversions should be76; necessary in the PTX.77define i32 @myctlz64_as_32(i64 %a) {78; CHECK-LABEL: myctlz64_as_32(79; CHECK: {80; CHECK-NEXT: .reg .b32 %r<2>;81; CHECK-NEXT: .reg .b64 %rd<2>;82; CHECK-EMPTY:83; CHECK-NEXT: // %bb.0:84; CHECK-NEXT: ld.param.b64 %rd1, [myctlz64_as_32_param_0];85; CHECK-NEXT: clz.b64 %r1, %rd1;86; CHECK-NEXT: st.param.b32 [func_retval0], %r1;87; CHECK-NEXT: ret;88 %val = call i64 @llvm.ctlz.i64(i64 %a, i1 false) readnone89 %trunc = trunc i64 %val to i3290 ret i32 %trunc91}92define i32 @myctlz64_as_32_2(i64 %a) {93; CHECK-LABEL: myctlz64_as_32_2(94; CHECK: {95; CHECK-NEXT: .reg .b32 %r<2>;96; CHECK-NEXT: .reg .b64 %rd<2>;97; CHECK-EMPTY:98; CHECK-NEXT: // %bb.0:99; CHECK-NEXT: ld.param.b64 %rd1, [myctlz64_as_32_2_param_0];100; CHECK-NEXT: clz.b64 %r1, %rd1;101; CHECK-NEXT: st.param.b32 [func_retval0], %r1;102; CHECK-NEXT: ret;103 %val = call i64 @llvm.ctlz.i64(i64 %a, i1 false) readnone104 %trunc = trunc i64 %val to i32105 ret i32 %trunc106}107 108; ctlz.i16 is implemented by extending the input to i32, computing the result,109; and then truncating the result back down to i16. But the NVPTX ABI110; zero-extends i16 return values to i32, so the final truncation doesn't appear111; in this function.112define i16 @myctlz_ret16(i16 %a) {113; CHECK-LABEL: myctlz_ret16(114; CHECK: {115; CHECK-NEXT: .reg .b32 %r<4>;116; CHECK-EMPTY:117; CHECK-NEXT: // %bb.0:118; CHECK-NEXT: ld.param.b16 %r1, [myctlz_ret16_param_0];119; CHECK-NEXT: clz.b32 %r2, %r1;120; CHECK-NEXT: add.s32 %r3, %r2, -16;121; CHECK-NEXT: st.param.b32 [func_retval0], %r3;122; CHECK-NEXT: ret;123 %val = call i16 @llvm.ctlz.i16(i16 %a, i1 false) readnone124 ret i16 %val125}126define i16 @myctlz_ret16_2(i16 %a) {127; CHECK-LABEL: myctlz_ret16_2(128; CHECK: {129; CHECK-NEXT: .reg .b32 %r<4>;130; CHECK-EMPTY:131; CHECK-NEXT: // %bb.0:132; CHECK-NEXT: ld.param.b16 %r1, [myctlz_ret16_2_param_0];133; CHECK-NEXT: shl.b32 %r2, %r1, 16;134; CHECK-NEXT: clz.b32 %r3, %r2;135; CHECK-NEXT: st.param.b32 [func_retval0], %r3;136; CHECK-NEXT: ret;137 %val = call i16 @llvm.ctlz.i16(i16 %a, i1 true) readnone138 ret i16 %val139}140 141; Here we store the result of ctlz.16 into an i16 pointer, so the trunc should142; remain.143define void @myctlz_store16(i16 %a, ptr %b) {144; CHECK-LABEL: myctlz_store16(145; CHECK: {146; CHECK-NEXT: .reg .b32 %r<4>;147; CHECK-NEXT: .reg .b64 %rd<2>;148; CHECK-EMPTY:149; CHECK-NEXT: // %bb.0:150; CHECK-NEXT: ld.param.b16 %r1, [myctlz_store16_param_0];151; CHECK-NEXT: clz.b32 %r2, %r1;152; CHECK-NEXT: add.s32 %r3, %r2, -16;153; CHECK-NEXT: ld.param.b64 %rd1, [myctlz_store16_param_1];154; CHECK-NEXT: st.b16 [%rd1], %r3;155; CHECK-NEXT: ret;156 %val = call i16 @llvm.ctlz.i16(i16 %a, i1 false) readnone157 store i16 %val, ptr %b158 ret void159}160define void @myctlz_store16_2(i16 %a, ptr %b) {161; CHECK-LABEL: myctlz_store16_2(162; CHECK: {163; CHECK-NEXT: .reg .b32 %r<4>;164; CHECK-NEXT: .reg .b64 %rd<2>;165; CHECK-EMPTY:166; CHECK-NEXT: // %bb.0:167; CHECK-NEXT: ld.param.b16 %r1, [myctlz_store16_2_param_0];168; CHECK-NEXT: clz.b32 %r2, %r1;169; CHECK-NEXT: add.s32 %r3, %r2, -16;170; CHECK-NEXT: ld.param.b64 %rd1, [myctlz_store16_2_param_1];171; CHECK-NEXT: st.b16 [%rd1], %r3;172; CHECK-NEXT: ret;173 %val = call i16 @llvm.ctlz.i16(i16 %a, i1 false) readnone174 store i16 %val, ptr %b175 ret void176}177