; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_30 -verify-machineinstrs | FileCheck %s
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_30 -verify-machineinstrs | %ptxas-verify %}


target triple = "nvptx-unknown-cuda"

declare { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64, i32)
declare i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1))

define ptx_kernel void @foo(i64 %img, ptr %red, i32 %idx) {
; CHECK-LABEL: foo(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<6>;
; CHECK-NEXT:    .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b64 %rd1, [foo_param_0];
; CHECK-NEXT:    ld.param.b64 %rd2, [foo_param_1];
; CHECK-NEXT:    cvta.to.global.u64 %rd3, %rd2;
; CHECK-NEXT:    ld.param.b32 %r1, [foo_param_2];
; CHECK-NEXT:    tex.1d.v4.f32.s32 {%r2, %r3, %r4, %r5}, [%rd1, {%r1}];
; CHECK-NEXT:    st.global.b32 [%rd3], %r2;
; CHECK-NEXT:    ret;
  %val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %img, i32 %idx)
  %ret = extractvalue { float, float, float, float } %val, 0
  store float %ret, ptr %red
  ret void
}


@tex0 = internal addrspace(1) global i64 0, align 8

define ptx_kernel void @bar(ptr %red, i32 %idx) {
; CHECK-LABEL: bar(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<6>;
; CHECK-NEXT:    .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b64 %rd1, [bar_param_0];
; CHECK-NEXT:    cvta.to.global.u64 %rd2, %rd1;
; CHECK-NEXT:    ld.param.b32 %r1, [bar_param_1];
; CHECK-NEXT:    tex.1d.v4.f32.s32 {%r2, %r3, %r4, %r5}, [tex0, {%r1}];
; CHECK-NEXT:    st.global.b32 [%rd2], %r2;
; CHECK-NEXT:    ret;
  %texHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @tex0)
  %val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %texHandle, i32 %idx)
  %ret = extractvalue { float, float, float, float } %val, 0
  store float %ret, ptr %red
  ret void
}

declare float @texfunc(i64)

define ptx_kernel void @baz(ptr %red, i32 %idx) {
; CHECK-LABEL: baz(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<9>;
; CHECK-NEXT:    .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b64 %rd1, [baz_param_0];
; CHECK-NEXT:    cvta.to.global.u64 %rd2, %rd1;
; CHECK-NEXT:    ld.param.b32 %r1, [baz_param_1];
; CHECK-NEXT:    mov.u64 %rd3, tex0;
; CHECK-NEXT:    tex.1d.v4.f32.s32 {%r2, %r3, %r4, %r5}, [tex0, {%r1}];
; CHECK-NEXT:    { // callseq 0, 0
; CHECK-NEXT:    .param .b64 param0;
; CHECK-NEXT:    .param .b32 retval0;
; CHECK-NEXT:    st.param.b64 [param0], %rd3;
; CHECK-NEXT:    call.uni (retval0), texfunc, (param0);
; CHECK-NEXT:    ld.param.b32 %r6, [retval0];
; CHECK-NEXT:    } // callseq 0
; CHECK-NEXT:    add.rn.f32 %r8, %r2, %r6;
; CHECK-NEXT:    st.global.b32 [%rd2], %r8;
; CHECK-NEXT:    ret;
  %texHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @tex0)
  %val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %texHandle, i32 %idx)
  %ret = extractvalue { float, float, float, float } %val, 0
  %texcall = tail call float @texfunc(i64 %texHandle)
  %ret2 = fadd float %ret, %texcall
  store float %ret2, ptr %red
  ret void
}

!nvvm.annotations = !{!1}
!1 = !{ptr addrspace(1) @tex0, !"texture", i32 1}