llvm/test/CodeGen/NVPTX/address-folder.ll

; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s | FileCheck %s
; RUN: llc -O0 < %s | FileCheck %s --check-prefix=O0

target triple = "nvptx64-nvidia-cuda"

; A byval kernel parameter loaded at a constant offset from a block other than
; the one that defines its address must still be read directly from parameter
; space (ld.param [param+off]), not via a materialized base register.
define ptx_kernel void @test_kernel_cross_block(ptr byval([2 x i64]) align 8 %r, i64 %n, ptr addrspace(1) %out) {
; CHECK-LABEL: test_kernel_cross_block(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b64 %rd<6>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b64 %rd2, [test_kernel_cross_block_param_1];
; CHECK-NEXT:    setp.lt.s64 %p1, %rd2, 1;
; CHECK-NEXT:    @%p1 bra $L__BB0_2;
; CHECK-NEXT:  // %bb.1: // %body
; CHECK-NEXT:    ld.param.b64 %rd1, [test_kernel_cross_block_param_2];
; CHECK-NEXT:    ld.param.b64 %rd3, [test_kernel_cross_block_param_0];
; CHECK-NEXT:    ld.param.b64 %rd4, [test_kernel_cross_block_param_0+8];
; CHECK-NEXT:    mul.lo.s64 %rd5, %rd3, %rd4;
; CHECK-NEXT:    st.global.b64 [%rd1], %rd5;
; CHECK-NEXT:  $L__BB0_2: // %done
; CHECK-NEXT:    ret;
;
; O0-LABEL: test_kernel_cross_block(
; O0:       {
; O0-NEXT:    .reg .pred %p<2>;
; O0-NEXT:    .reg .b64 %rd<7>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b64 %rd3, [test_kernel_cross_block_param_2];
; O0-NEXT:    ld.param.b64 %rd2, [test_kernel_cross_block_param_1];
; O0-NEXT:    mov.b64 %rd1, test_kernel_cross_block_param_0;
; O0-NEXT:    setp.lt.s64 %p1, %rd2, 1;
; O0-NEXT:    @%p1 bra $L__BB0_2;
; O0-NEXT:    bra.uni $L__BB0_1;
; O0-NEXT:  $L__BB0_1: // %body
; O0-NEXT:    ld.param.b64 %rd4, [%rd1];
; O0-NEXT:    ld.param.b64 %rd5, [%rd1+8];
; O0-NEXT:    mul.lo.s64 %rd6, %rd4, %rd5;
; O0-NEXT:    st.global.b64 [%rd3], %rd6;
; O0-NEXT:    bra.uni $L__BB0_2;
; O0-NEXT:  $L__BB0_2: // %done
; O0-NEXT:    ret;
  %c = icmp sgt i64 %n, 0
  br i1 %c, label %body, label %done
body:
  %v0 = load i64, ptr %r, align 8
  %p1 = getelementptr inbounds i8, ptr %r, i64 8
  %v1 = load i64, ptr %p1, align 8
  %m = mul i64 %v0, %v1
  store i64 %m, ptr addrspace(1) %out, align 8
  br label %done
done:
  ret void
}

; When the address is also used for arithmetic (here a variable offset) it is
; not just loaded, so the `mov` must be kept rather than folded into the load.
define ptx_kernel void @test_kernel_var_offset(ptr byval([8 x i64]) align 8 %r, i64 %n, ptr addrspace(1) %out) {
; CHECK-LABEL: test_kernel_var_offset(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b64 %rd<6>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b64 %rd4, [test_kernel_var_offset_param_1];
; CHECK-NEXT:    setp.lt.s64 %p1, %rd4, 1;
; CHECK-NEXT:    @%p1 bra $L__BB1_2;
; CHECK-NEXT:  // %bb.1: // %body
; CHECK-NEXT:    ld.param.b64 %rd2, [test_kernel_var_offset_param_2];
; CHECK-NEXT:    mov.b64 %rd3, test_kernel_var_offset_param_0;
; CHECK-NEXT:    add.s64 %rd1, %rd3, %rd4;
; CHECK-NEXT:    ld.param.b64 %rd5, [%rd1];
; CHECK-NEXT:    st.global.b64 [%rd2], %rd5;
; CHECK-NEXT:  $L__BB1_2: // %done
; CHECK-NEXT:    ret;
;
; O0-LABEL: test_kernel_var_offset(
; O0:       {
; O0-NEXT:    .reg .pred %p<2>;
; O0-NEXT:    .reg .b64 %rd<6>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b64 %rd3, [test_kernel_var_offset_param_2];
; O0-NEXT:    ld.param.b64 %rd2, [test_kernel_var_offset_param_1];
; O0-NEXT:    mov.b64 %rd1, test_kernel_var_offset_param_0;
; O0-NEXT:    setp.lt.s64 %p1, %rd2, 1;
; O0-NEXT:    @%p1 bra $L__BB1_2;
; O0-NEXT:    bra.uni $L__BB1_1;
; O0-NEXT:  $L__BB1_1: // %body
; O0-NEXT:    add.s64 %rd4, %rd1, %rd2;
; O0-NEXT:    ld.param.b64 %rd5, [%rd4];
; O0-NEXT:    st.global.b64 [%rd3], %rd5;
; O0-NEXT:    bra.uni $L__BB1_2;
; O0-NEXT:  $L__BB1_2: // %done
; O0-NEXT:    ret;
  %c = icmp sgt i64 %n, 0
  br i1 %c, label %body, label %done
body:
  %p = getelementptr inbounds i8, ptr %r, i64 %n
  %v = load i64, ptr %p, align 8
  store i64 %v, ptr addrspace(1) %out, align 8
  br label %done
done:
  ret void
}