48 lines · plain
1; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_50 -verify-machineinstrs | FileCheck %s2; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_50 -verify-machineinstrs | %ptxas-verify %}3 4; calls with a bitcasted function symbol should be fine, but in combination with5; a byval attribute were causing a segfault during isel. This testcase was6; reduced from a SYCL kernel using aggregate types which ended up being passed7; `byval`8 9target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"10target triple = "nvptx64-nvidia-cuda"11 12%"class.complex" = type { %"class.sycl::_V1::detail::half_impl::half", %"class.sycl::_V1::detail::half_impl::half" }13%"class.sycl::_V1::detail::half_impl::half" = type { half }14%complex_half = type { half, half }15 16; CHECK: .param .align 2 .b8 param2[4];17; CHECK: .param .align 2 .b8 retval0[4];18; CHECK-DAG: st.param.b16 [param2], %rs{{[0-9]+}};19; CHECK-DAG: st.param.b16 [param2+2], %rs{{[0-9]+}};20; CHECK: prototype_0 : .callprototype (.param .align 2 .b8 _[4]) _ (.param .b32 _, .param .b32 _, .param .align 2 .b8 _[4]);21; CHECK: call (retval0),22define weak_odr void @foo() {23entry:24 %call.i.i.i = tail call %"class.complex" @_Z20__spirv_GroupCMulKHRjjN5__spv12complex_halfE(i32 0, i32 0, ptr byval(%"class.complex") null)25 ret void26}27 28;; Function pointers can escape, so we have to use a conservative29;; alignment for a function that has address taken.30;;31declare ptr @usefp(ptr %fp)32; CHECK: .func callee(33; CHECK-NEXT: .param .align 2 .b8 callee_param_0[4]34define internal void @callee(ptr byval(%"class.complex") %byval_arg) {35 ret void36}37define void @boom() {38 %fp = call ptr @usefp(ptr @callee)39 ; CHECK-DAG: .param .align 2 .b8 param0[4];40 ; CHECK-DAG: st.param.b16 [param0], %rs{{[0-9]+}};41 ; CHECK-DAG: st.param.b16 [param0+2], %rs{{[0-9]+}};42 ; CHECK-DAG: .callprototype ()_ (.param .align 2 .b8 _[4]);43 call void %fp(ptr byval(%"class.complex") null)44 ret void45}46 47declare %complex_half @_Z20__spirv_GroupCMulKHRjjN5__spv12complex_halfE(i32, i32, ptr byval(%"class.complex"))48