diff options
Diffstat (limited to 'llvm/test/CodeGen/NVPTX/pr126337.ll')
| -rw-r--r-- | llvm/test/CodeGen/NVPTX/pr126337.ll | 41 |
1 files changed, 41 insertions, 0 deletions
diff --git a/llvm/test/CodeGen/NVPTX/pr126337.ll b/llvm/test/CodeGen/NVPTX/pr126337.ll new file mode 100644 index 0000000..95258f7 --- /dev/null +++ b/llvm/test/CodeGen/NVPTX/pr126337.ll @@ -0,0 +1,41 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 +; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_70 | FileCheck %s +; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_70 | %ptxas -arch=sm_70 -c - %} + +; This IR should compile without triggering assertions in LICM +; when the CopyToReg from %0 in the first BB gets eliminated +; but we still use its result in the second BB. +; Technically the problem happens in MIR, but there are multiple +; passes involved, so testing with the IR reproducer is more convenient. +; https://github.com/llvm/llvm-project/pull/126337#issuecomment-3081431594 + +target datalayout = "e-p6:32:32-i64:64-i128:128-v16:16-v32:32-n16:32:64" +target triple = "nvptx64-nvidia-cuda" + +define ptx_kernel void @Equal_GPU_DT_COMPLEX64_DT_BOOL_kernel(<2 x float> %0) { +; CHECK-LABEL: Equal_GPU_DT_COMPLEX64_DT_BOOL_kernel( +; CHECK: { +; CHECK-NEXT: .reg .pred %p<2>; +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<2>; +; CHECK-NEXT: .reg .b64 %rd<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: // %.preheader15 +; CHECK-NEXT: ld.param.b64 %rd1, [Equal_GPU_DT_COMPLEX64_DT_BOOL_kernel_param_0]; +; CHECK-NEXT: { .reg .b32 tmp; mov.b64 {%r1, tmp}, %rd1; } +; CHECK-NEXT: setp.eq.f32 %p1, %r1, 0f00000000; +; CHECK-NEXT: selp.b16 %rs1, 1, 0, %p1; +; CHECK-NEXT: $L__BB0_1: // =>This Inner Loop Header: Depth=1 +; CHECK-NEXT: mov.b64 %rd2, 0; +; CHECK-NEXT: st.b8 [%rd2], %rs1; +; CHECK-NEXT: bra.uni $L__BB0_1; +.preheader15: + br label %1 + +1: ; preds = %1, %.preheader15 + %2 = fcmp oeq <2 x float> %0, zeroinitializer + %3 = extractelement <2 x i1> %2, i64 0 + store i1 %3, ptr null, align 4 + br label %1 +} + |