63 lines · plain
1; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s2; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}3 4target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"5target triple = "nvptx64-nvidia-cuda"6 7; Verify that both %input and %output are converted to global pointers and then8; addrspacecast'ed back to the original type.9define ptx_kernel void @kernel(ptr %input, ptr %output) {10; CHECK-LABEL: .visible .entry kernel(11; CHECK: cvta.to.global.u6412; CHECK: cvta.to.global.u6413 %1 = load float, ptr %input, align 414; CHECK: ld.global.b3215 store float %1, ptr %output, align 416; CHECK: st.global.b3217 ret void18}19 20define ptx_kernel void @kernel2(ptr addrspace(1) %input, ptr addrspace(1) %output) {21; CHECK-LABEL: .visible .entry kernel2(22; CHECK-NOT: cvta.to.global.u6423 %1 = load float, ptr addrspace(1) %input, align 424; CHECK: ld.global.b3225 store float %1, ptr addrspace(1) %output, align 426; CHECK: st.global.b3227 ret void28}29 30%struct.S = type { ptr, ptr }31 32define ptx_kernel void @ptr_in_byval_kernel(ptr byval(%struct.S) %input, ptr %output) {33; CHECK-LABEL: .visible .entry ptr_in_byval_kernel(34; CHECK: ld.param.b64 %[[optr:rd.*]], [ptr_in_byval_kernel_param_1]35; CHECK: cvta.to.global.u64 %[[optr_g:.*]], %[[optr]];36; CHECK: ld.param.b64 %[[iptr:rd.*]], [ptr_in_byval_kernel_param_0+8]37; CHECK: cvta.to.global.u64 %[[iptr_g:.*]], %[[iptr]];38 %b_ptr = getelementptr inbounds %struct.S, ptr %input, i64 0, i32 139 %b = load ptr, ptr %b_ptr, align 840 %v = load i32, ptr %b, align 441; CHECK: ld.global.b32 %[[val:.*]], [%[[iptr_g]]]42 store i32 %v, ptr %output, align 443; CHECK: st.global.b32 [%[[optr_g]]], %[[val]]44 ret void45}46 47; Regular functions lower byval arguments differently. We need to make48; sure that we're loading byval argument data using [symbol+offset].49; There's also no assumption that all pointers within are in global space.50define void @ptr_in_byval_func(ptr byval(%struct.S) %input, ptr %output) {51; CHECK-LABEL: .visible .func ptr_in_byval_func(52; CHECK: ld.param.b64 %[[optr:rd.*]], [ptr_in_byval_func_param_1]53; CHECK: ld.param.b64 %[[iptr:rd.*]], [ptr_in_byval_func_param_0+8]54 %b_ptr = getelementptr inbounds %struct.S, ptr %input, i64 0, i32 155 %b = load ptr, ptr %b_ptr, align 856 %v = load i32, ptr %b, align 457; CHECK: ld.b32 %[[val:.*]], [%[[iptr]]]58 store i32 %v, ptr %output, align 459; CHECK: st.b32 [%[[optr]]], %[[val]]60 ret void61}62 63