1; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s 2; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} 3 4; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s 5; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %} 6 7; CHECK-DAG: %[[#EventTy:]] = OpTypeEvent 8; CHECK-DAG: %[[#StructEventTy:]] = OpTypeStruct %[[#EventTy]] 9; CHECK-DAG: %[[#FunPtrStructEventTy:]] = OpTypePointer Function %[[#StructEventTy]] 10; CHECK-DAG: %[[#GenPtrEventTy:]] = OpTypePointer Generic %[[#EventTy]] 11; CHECK-DAG: %[[#FunPtrEventTy:]] = OpTypePointer Function %[[#EventTy]] 12; CHECK: OpFunction 13; CHECK: %[[#Var:]] = OpVariable %[[#FunPtrStructEventTy]] Function 14; CHECK-NEXT: %[[#FunEvent:]] = OpBitcast %[[#FunPtrEventTy]] %[[#Var]] 15; CHECK-NEXT: %[[#GenEvent:]] = OpPtrCastToGeneric %[[#GenPtrEventTy]] %[[#FunEvent]] 16; CHECK-NEXT: OpGroupWaitEvents %[[#]] %[[#]] %[[#GenEvent]] 17 18%"class.sycl::_V1::device_event" = type { target("spirv.Event") } 19 20define weak_odr dso_local spir_kernel void @foo() { 21entry: 22 %var = alloca %"class.sycl::_V1::device_event" 23 %eventptr = addrspacecast ptr %var to ptr addrspace(4) 24 call spir_func void @_Z23__spirv_GroupWaitEventsjiP9ocl_event(i32 2, i32 1, ptr addrspace(4) %eventptr) 25 ret void 26} 27 28declare dso_local spir_func void @_Z23__spirv_GroupWaitEventsjiP9ocl_event(i32, i32, ptr addrspace(4)) 29