; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_75 -mattr=+ptx82 | FileCheck %s --check-prefix=PTX82
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_75 -mattr=+ptx83 | FileCheck %s --check-prefix=PTX83
; RUN: %if ptxas-sm_75 && ptxas-isa-8.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_75 -mattr=+ptx82 | %ptxas-verify -arch=sm_75 %}
; RUN: %if ptxas-sm_75 && ptxas-isa-8.3 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_75 -mattr=+ptx83 | %ptxas-verify -arch=sm_75 %}

target triple = "nvptx64-nvidia-cuda"

define internal i32 @callee(i32 %x) noinline {
; PTX82-LABEL: callee(
; PTX82:       {
; PTX82-NEXT:    .reg .b32 %r<3>;
; PTX82-EMPTY:
; PTX82-NEXT:  // %bb.0:
; PTX82-NEXT:    ld.param.b32 %r1, [callee_param_0];
; PTX82-NEXT:    add.s32 %r2, %r1, 1;
; PTX82-NEXT:    st.param.b32 [func_retval0], %r2;
; PTX82-NEXT:    ret;
;
; PTX83-LABEL: callee(
; PTX83:       {
; PTX83-NEXT:    .reg .b32 %r<3>;
; PTX83-EMPTY:
; PTX83-NEXT:  // %bb.0:
; PTX83-NEXT:    ld.param::func.b32 %r1, [callee_param_0];
; PTX83-NEXT:    add.s32 %r2, %r1, 1;
; PTX83-NEXT:    st.param::func.b32 [func_retval0], %r2;
; PTX83-NEXT:    ret;
  %y = add i32 %x, 1
  ret i32 %y
}

define ptx_kernel void @kernel(i32 %x, ptr %out) {
; PTX82-LABEL: kernel(
; PTX82:       {
; PTX82-NEXT:    .reg .b32 %r<3>;
; PTX82-NEXT:    .reg .b64 %rd<3>;
; PTX82-EMPTY:
; PTX82-NEXT:  // %bb.0:
; PTX82-NEXT:    ld.param.b32 %r1, [kernel_param_0];
; PTX82-NEXT:    ld.param.b64 %rd1, [kernel_param_1];
; PTX82-NEXT:    cvta.to.global.u64 %rd2, %rd1;
; PTX82-NEXT:    { // callseq 0, 0
; PTX82-NEXT:    .param .b32 param0;
; PTX82-NEXT:    .param .b32 retval0;
; PTX82-NEXT:    st.param.b32 [param0], %r1;
; PTX82-NEXT:    call.uni (retval0), callee, (param0);
; PTX82-NEXT:    ld.param.b32 %r2, [retval0];
; PTX82-NEXT:    } // callseq 0
; PTX82-NEXT:    st.global.b32 [%rd2], %r2;
; PTX82-NEXT:    ret;
;
; PTX83-LABEL: kernel(
; PTX83:       {
; PTX83-NEXT:    .reg .b32 %r<3>;
; PTX83-NEXT:    .reg .b64 %rd<3>;
; PTX83-EMPTY:
; PTX83-NEXT:  // %bb.0:
; PTX83-NEXT:    ld.param::entry.b32 %r1, [kernel_param_0];
; PTX83-NEXT:    ld.param::entry.b64 %rd1, [kernel_param_1];
; PTX83-NEXT:    cvta.to.global.u64 %rd2, %rd1;
; PTX83-NEXT:    { // callseq 0, 0
; PTX83-NEXT:    .param .b32 param0;
; PTX83-NEXT:    .param .b32 retval0;
; PTX83-NEXT:    st.param::func.b32 [param0], %r1;
; PTX83-NEXT:    call.uni (retval0), callee, (param0);
; PTX83-NEXT:    ld.param::func.b32 %r2, [retval0];
; PTX83-NEXT:    } // callseq 0
; PTX83-NEXT:    st.global.b32 [%rd2], %r2;
; PTX83-NEXT:    ret;
  %ret = call i32 @callee(i32 %x)
  store i32 %ret, ptr %out
  ret void
}

; A device function that loads through a ptr addrspace(101) (entry param ref
; passed from a kernel). The ::entry subqualifier on the register-indirect
; ld.param is required and older ptxas miscompiles this without it.

define internal i32 @callee_entry_param(ptr addrspace(101) %p) noinline {
; PTX82-LABEL: callee_entry_param(
; PTX82:       {
; PTX82-NEXT:    .reg .b32 %r<2>;
; PTX82-NEXT:    .reg .b64 %rd<2>;
; PTX82-EMPTY:
; PTX82-NEXT:  // %bb.0:
; PTX82-NEXT:    ld.param.b64 %rd1, [callee_entry_param_param_0];
; PTX82-NEXT:    ld.param.b32 %r1, [%rd1];
; PTX82-NEXT:    st.param.b32 [func_retval0], %r1;
; PTX82-NEXT:    ret;
;
; PTX83-LABEL: callee_entry_param(
; PTX83:       {
; PTX83-NEXT:    .reg .b32 %r<2>;
; PTX83-NEXT:    .reg .b64 %rd<2>;
; PTX83-EMPTY:
; PTX83-NEXT:  // %bb.0:
; PTX83-NEXT:    ld.param::func.b64 %rd1, [callee_entry_param_param_0];
; PTX83-NEXT:    ld.param::entry.b32 %r1, [%rd1];
; PTX83-NEXT:    st.param::func.b32 [func_retval0], %r1;
; PTX83-NEXT:    ret;
  %v = load i32, ptr addrspace(101) %p, align 4
  ret i32 %v
}

define ptx_kernel void @kernel_pass_param_ref(ptr byval(i32) align 4 "nvvm.grid_constant" %input, ptr %out) {
; PTX82-LABEL: kernel_pass_param_ref(
; PTX82:       {
; PTX82-NEXT:    .reg .b32 %r<2>;
; PTX82-NEXT:    .reg .b64 %rd<4>;
; PTX82-EMPTY:
; PTX82-NEXT:  // %bb.0:
; PTX82-NEXT:    mov.b64 %rd1, kernel_pass_param_ref_param_0;
; PTX82-NEXT:    ld.param.b64 %rd2, [kernel_pass_param_ref_param_1];
; PTX82-NEXT:    cvta.to.global.u64 %rd3, %rd2;
; PTX82-NEXT:    { // callseq 1, 0
; PTX82-NEXT:    .param .b64 param0;
; PTX82-NEXT:    .param .b32 retval0;
; PTX82-NEXT:    st.param.b64 [param0], %rd1;
; PTX82-NEXT:    call.uni (retval0), callee_entry_param, (param0);
; PTX82-NEXT:    ld.param.b32 %r1, [retval0];
; PTX82-NEXT:    } // callseq 1
; PTX82-NEXT:    st.global.b32 [%rd3], %r1;
; PTX82-NEXT:    ret;
;
; PTX83-LABEL: kernel_pass_param_ref(
; PTX83:       {
; PTX83-NEXT:    .reg .b32 %r<2>;
; PTX83-NEXT:    .reg .b64 %rd<4>;
; PTX83-EMPTY:
; PTX83-NEXT:  // %bb.0:
; PTX83-NEXT:    mov.b64 %rd1, kernel_pass_param_ref_param_0;
; PTX83-NEXT:    ld.param::entry.b64 %rd2, [kernel_pass_param_ref_param_1];
; PTX83-NEXT:    cvta.to.global.u64 %rd3, %rd2;
; PTX83-NEXT:    { // callseq 1, 0
; PTX83-NEXT:    .param .b64 param0;
; PTX83-NEXT:    .param .b32 retval0;
; PTX83-NEXT:    st.param::func.b64 [param0], %rd1;
; PTX83-NEXT:    call.uni (retval0), callee_entry_param, (param0);
; PTX83-NEXT:    ld.param::func.b32 %r1, [retval0];
; PTX83-NEXT:    } // callseq 1
; PTX83-NEXT:    st.global.b32 [%rd3], %r1;
; PTX83-NEXT:    ret;
  %p = addrspacecast ptr %input to ptr addrspace(101)
  %ret = call i32 @callee_entry_param(ptr addrspace(101) %p)
  store i32 %ret, ptr %out
  ret void
}