41 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_70 | FileCheck %s3; RUN: %if ptxas-sm_70 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_70 | %ptxas-verify -arch=sm_70 %}4 5; This IR should compile without triggering assertions in LICM6; when the CopyToReg from %0 in the first BB gets eliminated7; but we still use its result in the second BB.8; Technically the problem happens in MIR, but there are multiple9; passes involved, so testing with the IR reproducer is more convenient.10; https://github.com/llvm/llvm-project/pull/126337#issuecomment-308143159411 12target datalayout = "e-p6:32:32-i64:64-i128:128-v16:16-v32:32-n16:32:64"13target triple = "nvptx64-nvidia-cuda"14 15define ptx_kernel void @Equal_GPU_DT_COMPLEX64_DT_BOOL_kernel(<2 x float> %0) {16; CHECK-LABEL: Equal_GPU_DT_COMPLEX64_DT_BOOL_kernel(17; CHECK: {18; CHECK-NEXT: .reg .pred %p<2>;19; CHECK-NEXT: .reg .b16 %rs<2>;20; CHECK-NEXT: .reg .b32 %r<3>;21; CHECK-NEXT: .reg .b64 %rd<2>;22; CHECK-EMPTY:23; CHECK-NEXT: // %bb.0: // %.preheader1524; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [Equal_GPU_DT_COMPLEX64_DT_BOOL_kernel_param_0];25; CHECK-NEXT: setp.eq.f32 %p1, %r1, 0f00000000;26; CHECK-NEXT: selp.b16 %rs1, 1, 0, %p1;27; CHECK-NEXT: $L__BB0_1: // =>This Inner Loop Header: Depth=128; CHECK-NEXT: mov.b64 %rd1, 0;29; CHECK-NEXT: st.b8 [%rd1], %rs1;30; CHECK-NEXT: bra.uni $L__BB0_1;31.preheader15:32 br label %133 341: ; preds = %1, %.preheader1535 %2 = fcmp oeq <2 x float> %0, zeroinitializer36 %3 = extractelement <2 x i1> %2, i64 037 store i1 %3, ptr null, align 438 br label %139}40 41