59 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_30 -verify-machineinstrs | FileCheck %s4; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}5; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_30 -verify-machineinstrs | %ptxas-verify %}6 7target triple = "nvptx-unknown-cuda"8 9declare i32 @llvm.nvvm.suld.1d.i32.trap(i64, i32)10declare i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1))11 12 13define ptx_kernel void @foo(i64 %img, ptr %red, i32 %idx) {14; CHECK-LABEL: foo(15; CHECK: {16; CHECK-NEXT: .reg .b32 %r<4>;17; CHECK-NEXT: .reg .b64 %rd<4>;18; CHECK-EMPTY:19; CHECK-NEXT: // %bb.0:20; CHECK-NEXT: ld.param.b64 %rd1, [foo_param_0];21; CHECK-NEXT: ld.param.b64 %rd2, [foo_param_1];22; CHECK-NEXT: cvta.to.global.u64 %rd3, %rd2;23; CHECK-NEXT: ld.param.b32 %r1, [foo_param_2];24; CHECK-NEXT: suld.b.1d.b32.trap {%r2}, [%rd1, {%r1}];25; CHECK-NEXT: cvt.rn.f32.s32 %r3, %r2;26; CHECK-NEXT: st.global.b32 [%rd3], %r3;27; CHECK-NEXT: ret;28 %val = tail call i32 @llvm.nvvm.suld.1d.i32.trap(i64 %img, i32 %idx)29 %ret = sitofp i32 %val to float30 store float %ret, ptr %red31 ret void32}33 34@surf0 = internal addrspace(1) global i64 0, align 835 36define ptx_kernel void @bar(ptr %red, i32 %idx) {37; CHECK-LABEL: bar(38; CHECK: {39; CHECK-NEXT: .reg .b32 %r<4>;40; CHECK-NEXT: .reg .b64 %rd<3>;41; CHECK-EMPTY:42; CHECK-NEXT: // %bb.0:43; CHECK-NEXT: ld.param.b64 %rd1, [bar_param_0];44; CHECK-NEXT: cvta.to.global.u64 %rd2, %rd1;45; CHECK-NEXT: ld.param.b32 %r1, [bar_param_1];46; CHECK-NEXT: suld.b.1d.b32.trap {%r2}, [surf0, {%r1}];47; CHECK-NEXT: cvt.rn.f32.s32 %r3, %r2;48; CHECK-NEXT: st.global.b32 [%rd2], %r3;49; CHECK-NEXT: ret;50 %surfHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @surf0)51 %val = tail call i32 @llvm.nvvm.suld.1d.i32.trap(i64 %surfHandle, i32 %idx)52 %ret = sitofp i32 %val to float53 store float %ret, ptr %red54 ret void55}56 57!nvvm.annotations = !{!1}58!1 = !{ptr addrspace(1) @surf0, !"surface", i32 1}59