| ; 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 |
| } |
| |
