36 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mcpu=sm_20 -O0 | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mcpu=sm_20 -O0 | %ptxas-verify %}4 5target triple = "nvptx64-unknown-unknown"6 7define ptr @test1(ptr %p) {8; CHECK-LABEL: test1(9; CHECK: {10; CHECK-NEXT: .reg .b64 %rd<2>;11; CHECK-EMPTY:12; CHECK-NEXT: // %bb.0:13; CHECK-NEXT: ld.param.b64 %rd1, [test1_param_0];14; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;15; CHECK-NEXT: ret;16 %a = addrspacecast ptr %p to ptr addrspace(5)17 %b = addrspacecast ptr addrspace(5) %a to ptr18 ret ptr %b19}20 21define ptr addrspace(1) @test2(ptr addrspace(5) %p) {22; CHECK-LABEL: test2(23; CHECK: {24; CHECK-NEXT: .reg .b64 %rd<4>;25; CHECK-EMPTY:26; CHECK-NEXT: // %bb.0:27; CHECK-NEXT: ld.param.b64 %rd1, [test2_param_0];28; CHECK-NEXT: cvta.local.u64 %rd2, %rd1;29; CHECK-NEXT: cvta.to.global.u64 %rd3, %rd2;30; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;31; CHECK-NEXT: ret;32 %a = addrspacecast ptr addrspace(5) %p to ptr33 %b = addrspacecast ptr %a to ptr addrspace(1)34 ret ptr addrspace(1) %b35}36