llvm/test/CodeGen/SPIRV/event-wait-ptr-type.ll

   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
  20 define weak_odr dso_local spir_kernel void @foo() {
  21 entry:
  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
  28 declare dso_local spir_func void @_Z23__spirv_GroupWaitEventsjiP9ocl_event(i32, i32, ptr addrspace(4))