[llvm] e6ab63e - [SPIRV] Fix enqueue empty kernel (#187671)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Jun 1 07:44:52 PDT 2026
Author: idubinov
Date: 2026-06-01T16:44:46+02:00
New Revision: e6ab63e40b5d53ec7c2b03cc6d084c0edeb7f32b
URL: https://github.com/llvm/llvm-project/commit/e6ab63e40b5d53ec7c2b03cc6d084c0edeb7f32b
DIFF: https://github.com/llvm/llvm-project/commit/e6ab63e40b5d53ec7c2b03cc6d084c0edeb7f32b.diff
LOG: [SPIRV] Fix enqueue empty kernel (#187671)
Function reference arguments don't get spv_bitcast after opaque pointer
migration, while data pointer arguments still might. Therefore:
a. getBlockStructInstr() got updated to check G_GLOBAL_VALUE ->
G_ADDRSPACE_CAST pattern for function reference arguments;
b. buildEnqueueKernel() got updated to add bitcast the block literal
pointer from struct* to i8*, as required by OpEnqueueKernel.
---------
Co-authored-by: Arseniy Obolenskiy <gooddoog at student.su>
Co-authored-by: Marcos Maronas <mmaronas at amd.com>
Co-authored-by: Dmitry Sidorov <dsidorov at amd.com>
Added:
Modified:
llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
llvm/test/CodeGen/SPIRV/transcoding/enqueue_kernel.ll
Removed:
llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll
################################################################################
diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
index 1febfc21ed3b0..f2607ffe0280b 100644
--- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
@@ -370,21 +370,18 @@ lookupBuiltin(StringRef DemangledCall,
static MachineInstr *getBlockStructInstr(Register ParamReg,
MachineRegisterInfo *MRI) {
- // We expect the following sequence of instructions:
- // %0:_(pN) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.alloca)
- // or = G_GLOBAL_VALUE @block_literal_global
- // %1:_(pN) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.bitcast), %0
- // %2:_(p4) = G_ADDRSPACE_CAST %1:_(pN)
+ // We expect ParamReg to be defined by G_ADDRSPACE_CAST with a source from
+ // G_GLOBAL_VALUE or spv_alloca. Returns the source instruction.
MachineInstr *MI = MRI->getUniqueVRegDef(ParamReg);
assert(MI->getOpcode() == TargetOpcode::G_ADDRSPACE_CAST &&
MI->getOperand(1).isReg());
Register BitcastReg = MI->getOperand(1).getReg();
MachineInstr *BitcastMI = MRI->getUniqueVRegDef(BitcastReg);
- assert(isSpvIntrinsic(*BitcastMI, Intrinsic::spv_bitcast) &&
- BitcastMI->getOperand(2).isReg());
- Register ValueReg = BitcastMI->getOperand(2).getReg();
- MachineInstr *ValueMI = MRI->getUniqueVRegDef(ValueReg);
- return ValueMI;
+ assert(BitcastMI && "Definition for source reg not found.");
+ if (BitcastMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
+ isSpvIntrinsic(*BitcastMI, Intrinsic::spv_alloca))
+ return BitcastMI;
+ llvm_unreachable("getBlockStructInstr: unexpected instruction pattern");
}
// Return an integer constant corresponding to the given register and
@@ -425,7 +422,7 @@ static const Type *getBlockStructType(Register ParamReg,
// section 6.12.5 should guarantee that we can do this.
MachineInstr *MI = getBlockStructInstr(ParamReg, MRI);
if (MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
- return MI->getOperand(1).getGlobal()->getType();
+ return MI->getOperand(1).getGlobal()->getValueType();
assert(isSpvIntrinsic(*MI, Intrinsic::spv_alloca) &&
"Blocks in OpenCL C must be traceable to allocation site");
return getMachineInstrType(MI);
@@ -2924,99 +2921,159 @@ static bool buildNDRange(const SPIRV::IncomingCall *Call,
.addUse(TmpReg);
}
-// TODO: maybe move to the global register.
-static SPIRVTypeInst
-getOrCreateSPIRVDeviceEventPointer(MachineIRBuilder &MIRBuilder,
- SPIRVGlobalRegistry *GR) {
- LLVMContext &Context = MIRBuilder.getMF().getFunction().getContext();
- unsigned SC1 = storageClassToAddressSpace(SPIRV::StorageClass::Generic);
- Type *PtrType = PointerType::get(Context, SC1);
- return GR->getOrCreateSPIRVType(PtrType, MIRBuilder,
- SPIRV::AccessQualifier::ReadWrite, true);
-}
-
static bool buildEnqueueKernel(const SPIRV::IncomingCall *Call,
MachineIRBuilder &MIRBuilder,
SPIRVGlobalRegistry *GR) {
+ // In this function there are three stages:
+ // 1. prepare call indexes in order we expect them.
+ // 2. process all arguments which requered preparation.
+ // 3. create a SPIRV operator with arguments.
+
MachineRegisterInfo *MRI = MIRBuilder.getMRI();
const DataLayout &DL = MIRBuilder.getDataLayout();
- bool IsSpirvOp = Call->isSpirvOp();
- bool HasEvents = Call->Builtin->Name.contains("events") || IsSpirvOp;
const SPIRVTypeInst Int32Ty = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
- // Make vararg instructions before OpEnqueueKernel.
- // Local sizes arguments: Sizes of block invoke arguments. Clang generates
- // local size operands as an array, so we need to unpack them.
+ // 1. prepare call indexes in order we expect them.
+ // Based on clang sources, clang/lib/CodeGen/CGBuiltin.cpp, BIenqueue_kernel,
+ // We expect 4
diff erent layouts of call arguments:
+ // 1) No events, no vargs: {Queue, Flags, Range, Kernel, Block};
+ // 2) No events, varargs: {Queue, Flags, Range, Kernel, Block, NumElem,
+ // ElemPtr};
+ // 3) events, no varargs: {Queue, Flags, Range, NumEvents,
+ // EventWaitList, EventRet, Kernel, Block};
+ // 4) events, varargs: {Queue,
+ // Flags, Range, NumEvents, EventWaitList, EventRet, Kernel, Block,
+ // NumElem, ElemPtr};
+ //
+ // We also may expect __spirv_EnqueueKernel
+
+ bool IsSpirvOp = Call->isSpirvOp();
+ bool HasEvents = Call->Builtin->Name.contains("_events") || IsSpirvOp;
+ bool HasVarArgs = Call->Builtin->Name.contains("_varargs") || IsSpirvOp;
+
+ const unsigned NumArgs = Call->Arguments.size();
+ const unsigned BaseArgIdx = 0;
+ const unsigned IncorrectIdx = NumArgs + 1;
+
+ const unsigned QueueIdx = BaseArgIdx;
+ const unsigned FlagsIdx = BaseArgIdx + 1;
+ const unsigned NDRangeIdx = BaseArgIdx + 2;
+ const unsigned NumEventsIdx = HasEvents ? BaseArgIdx + 3 : IncorrectIdx;
+ const unsigned WaitEventsIdx = HasEvents ? BaseArgIdx + 4 : IncorrectIdx;
+ const unsigned RetEventIdx = HasEvents ? BaseArgIdx + 5 : IncorrectIdx;
+ const unsigned InvokeIdx = BaseArgIdx + 3 + (HasEvents ? 3 : 0);
+ const unsigned ParamIdx = BaseArgIdx + 4 + (HasEvents ? 3 : 0);
+ const unsigned LocalSizeNumElemIdx =
+ HasVarArgs ? (BaseArgIdx + 5 + (HasEvents ? 3 : 0)) : IncorrectIdx;
+ const unsigned LocalSizeElemPtrIdx =
+ HasVarArgs ? (BaseArgIdx + 6 + (HasEvents ? 3 : 0)) : IncorrectIdx;
+
+ const unsigned LastArgIdx =
+ (BaseArgIdx + 4 + (HasEvents ? 3 : 0) + (HasVarArgs ? 2 : 0));
+ assert(LastArgIdx < NumArgs && "Incorrect number arguments");
+
+ // 2. Process all arguments which requered preparation.
+ // 2.1 Events - use Call arguments, or use dummy nulls in case of absence of
+ // events
+ Register NumEventsReg;
+ Register WaitEventsReg;
+ Register RetEventReg;
+ if (HasEvents) {
+ NumEventsReg = Call->Arguments[NumEventsIdx];
+ WaitEventsReg = Call->Arguments[WaitEventsIdx];
+ RetEventReg = Call->Arguments[RetEventIdx];
+ } else {
+ NumEventsReg = buildConstantIntReg32(0, MIRBuilder, GR);
+ // Per SPIR-V spec, OpEnqueueKernel's Wait Events / Ret Event operands
+ // must be pointers to OpTypeDeviceEvent. Build the LLVM TargetExtType
+ // for spirv.DeviceEvent so getOrCreateSPIRVPointerType can register
+ // both the SPIR-V type and the LLVM-side mapping; deriving the pointer
+ // from an opaque LLVM PointerType would lose the pointee and deduce
+ // to <Generic i8*>.
+ LLVMContext &Ctx = MIRBuilder.getMF().getFunction().getContext();
+ Type *DeviceEventTy = TargetExtType::get(Ctx, "spirv.DeviceEvent");
+ SPIRVTypeInst DeviceEventPtrTy = GR->getOrCreateSPIRVPointerType(
+ DeviceEventTy, MIRBuilder, SPIRV::StorageClass::Generic);
+ Register NullPtr =
+ GR->getOrCreateConstNullPtr(MIRBuilder, DeviceEventPtrTy);
+ WaitEventsReg = NullPtr;
+ RetEventReg = NullPtr;
+ }
+
+ // 2.2 Invoke (Kernel)
+ // The Invoke operand of OpEnqueueKernel must be the function's <id>
+ // (per SPIR-V spec). The frontend hands us the result of an
+ // addrspacecast of @block_invoke_kernel; bypass that cast so the
+ // operand references the underlying G_GLOBAL_VALUE register, which
+ // selectGlobalValue lowers to a placeholder later rewritten by
+ // SPIRVModuleAnalysis to the OpFunction <id>.
+ MachineInstr *InvokeGlobalMI =
+ getBlockStructInstr(Call->Arguments[InvokeIdx], MRI);
+ assert(InvokeGlobalMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE);
+ Register InvokeReg = InvokeGlobalMI->getOperand(0).getReg();
+ // OpEnqueueKernel's Invoke operand uses the pID register class.
+ MRI->setRegClass(InvokeReg, &SPIRV::pIDRegClass);
+
+ // 2.3 Param, Param Size, Param Align
+ Register BlockLiteralReg = Call->Arguments[ParamIdx];
+ const SPIRVTypeInst Int8Ty = GR->getOrCreateSPIRVIntegerType(8, MIRBuilder);
+ const SPIRVTypeInst Int8PtrGen = GR->getOrCreateSPIRVPointerType(
+ Int8Ty, MIRBuilder, SPIRV::StorageClass::Generic);
+ Type *PType = const_cast<Type *>(getBlockStructType(BlockLiteralReg, MRI));
+
+ Register ParamReg = createVirtualRegister(Int8PtrGen, GR, MIRBuilder);
+ MIRBuilder.buildInstr(SPIRV::OpBitcast)
+ .addDef(ParamReg)
+ .addUse(GR->getSPIRVTypeID(Int8PtrGen))
+ .addUse(BlockLiteralReg);
+ // TODO: these numbers should be obtained from block literal structure.
+ Register ParamSizeReg =
+ buildConstantIntReg32(DL.getTypeStoreSize(PType), MIRBuilder, GR);
+ Register ParamAlignReg =
+ buildConstantIntReg32(DL.getPrefTypeAlign(PType).value(), MIRBuilder, GR);
+
+ // 2.4 Local Size Array
SmallVector<Register, 16> LocalSizes;
- if (Call->Builtin->Name.contains("_varargs") || IsSpirvOp) {
- const unsigned LocalSizeArrayIdx = HasEvents ? 9 : 6;
- Register GepReg = Call->Arguments[LocalSizeArrayIdx];
- MachineInstr *GepMI = MRI->getUniqueVRegDef(GepReg);
- assert(isSpvIntrinsic(*GepMI, Intrinsic::spv_gep) &&
- GepMI->getOperand(3).isReg());
- Register ArrayReg = GepMI->getOperand(3).getReg();
- MachineInstr *ArrayMI = MRI->getUniqueVRegDef(ArrayReg);
- const Type *LocalSizeTy = getMachineInstrType(ArrayMI);
- assert(LocalSizeTy && "Local size type is expected");
- const uint64_t LocalSizeNum =
- cast<ArrayType>(LocalSizeTy)->getNumElements();
- unsigned SC = storageClassToAddressSpace(SPIRV::StorageClass::Generic);
- const LLT LLType = LLT::pointer(SC, GR->getPointerSize());
- const SPIRVTypeInst PointerSizeTy = GR->getOrCreateSPIRVPointerType(
- Int32Ty, MIRBuilder, SPIRV::StorageClass::Function);
- for (unsigned I = 0; I < LocalSizeNum; ++I) {
+ if (HasVarArgs) {
+ Register LocalSizeNumElem = Call->Arguments[LocalSizeNumElemIdx];
+ MachineInstr *LocalSizeNumElemMI = MRI->getUniqueVRegDef(LocalSizeNumElem);
+ const MachineOperand &ConstOp = LocalSizeNumElemMI->getOperand(1);
+ assert(LocalSizeNumElemMI->getOpcode() == TargetOpcode::G_CONSTANT &&
+ ConstOp.isCImm() && "Expected constant immediate");
+ uint64_t NumElem = ConstOp.getCImm()->getValue().getZExtValue();
+
+ Register LocalSizeArrayReg = Call->Arguments[LocalSizeElemPtrIdx];
+
+ for (unsigned i = 0; i < NumElem; ++i) {
Register Reg = MRI->createVirtualRegister(&SPIRV::pIDRegClass);
- MRI->setType(Reg, LLType);
- GR->assignSPIRVTypeToVReg(PointerSizeTy, Reg, MIRBuilder.getMF());
auto GEPInst = MIRBuilder.buildIntrinsic(
Intrinsic::spv_gep, ArrayRef<Register>{Reg}, true, false);
GEPInst
- .addImm(GepMI->getOperand(2).getImm()) // In bound.
- .addUse(ArrayMI->getOperand(0).getReg()) // Alloca.
+ .addImm(0) // In bound.
+ .addUse(LocalSizeArrayReg) // Base pointer.
.addUse(buildConstantIntReg32(0, MIRBuilder, GR)) // Indices.
- .addUse(buildConstantIntReg32(I, MIRBuilder, GR));
+ .addUse(buildConstantIntReg32(i, MIRBuilder, GR));
LocalSizes.push_back(Reg);
}
}
- // SPIRV OpEnqueueKernel instruction has 10+ arguments.
+ // 3. create a SPIRV operator with arguments.
auto MIB = MIRBuilder.buildInstr(SPIRV::OpEnqueueKernel)
.addDef(Call->ReturnRegister)
- .addUse(GR->getSPIRVTypeID(Int32Ty));
-
- // Copy all arguments before block invoke function pointer.
- const unsigned BlockFIdx = HasEvents ? 6 : 3;
- for (unsigned i = 0; i < BlockFIdx; i++)
- MIB.addUse(Call->Arguments[i]);
+ .addUse(GR->getSPIRVTypeID(Int32Ty))
+ .addUse(Call->Arguments[QueueIdx])
+ .addUse(Call->Arguments[FlagsIdx])
+ .addUse(Call->Arguments[NDRangeIdx])
+ .addUse(NumEventsReg)
+ .addUse(WaitEventsReg)
+ .addUse(RetEventReg)
+ .addUse(InvokeReg)
+ .addUse(ParamReg)
+ .addUse(ParamSizeReg)
+ .addUse(ParamAlignReg);
+ for (auto &LocalSize : LocalSizes)
+ MIB.addUse(LocalSize);
- // If there are no event arguments in the original call, add dummy ones.
- if (!HasEvents) {
- MIB.addUse(buildConstantIntReg32(0, MIRBuilder, GR)); // Dummy num events.
- Register NullPtr = GR->getOrCreateConstNullPtr(
- MIRBuilder, getOrCreateSPIRVDeviceEventPointer(MIRBuilder, GR));
- MIB.addUse(NullPtr); // Dummy wait events.
- MIB.addUse(NullPtr); // Dummy ret event.
- }
-
- MachineInstr *BlockMI = getBlockStructInstr(Call->Arguments[BlockFIdx], MRI);
- assert(BlockMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE);
- // Invoke: Pointer to invoke function.
- MIB.addGlobalAddress(BlockMI->getOperand(1).getGlobal());
-
- Register BlockLiteralReg = Call->Arguments[BlockFIdx + 1];
- // Param: Pointer to block literal.
- MIB.addUse(BlockLiteralReg);
-
- Type *PType = const_cast<Type *>(getBlockStructType(BlockLiteralReg, MRI));
- // TODO: these numbers should be obtained from block literal structure.
- // Param Size: Size of block literal structure.
- MIB.addUse(buildConstantIntReg32(DL.getTypeStoreSize(PType), MIRBuilder, GR));
- // Param Aligment: Aligment of block literal structure.
- MIB.addUse(buildConstantIntReg32(DL.getPrefTypeAlign(PType).value(),
- MIRBuilder, GR));
-
- for (unsigned i = 0; i < LocalSizes.size(); i++)
- MIB.addUse(LocalSizes[i]);
return true;
}
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index c575932cd8709..c3f21fe025bd5 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -6688,10 +6688,12 @@ bool SPIRVInstructionSelector::selectGlobalValue(
return true;
}
MachineInstrBuilder MIB3 =
- BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantNull))
+ BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpUndef))
.addDef(ResVReg)
.addUse(GR.getSPIRVTypeID(ResType));
GR.add(ConstVal, MIB3);
+ GR.recordFunctionPointer(&MIB3.getInstr()->getOperand(0),
+ cast<Function>(GV));
MIB3.constrainAllUses(TII, TRI, RBI);
return true;
}
diff --git a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
index 217e7619a9a3c..49d9dc95603ca 100644
--- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
@@ -341,7 +341,23 @@ bool SPIRVModuleAnalysis::isDeclSection(const MachineRegisterInfo &MRI,
return true;
}
if (GR->hasConstFunPtr() && Opcode == SPIRV::OpUndef) {
+ // The OpUndef may be a placeholder for a function reference recorded by
+ // selectGlobalValue. Skip emitting it if any user consumes it as a
+ // function-pointer-like operand (OpConstantFunctionPointerINTEL operand 2,
+ // or OpEnqueueKernel's Invoke operand at index 8). The rewrite happens
+ // in visitFunPtrUse, which aliases the OpUndef's vreg to the function's
+ // global <id>.
Register DefReg = MI.getOperand(0).getReg();
+ if (GR->getFunctionDefinitionByUse(&MI.getOperand(0))) {
+ for (MachineInstr &UseMI : MRI.use_instructions(DefReg)) {
+ unsigned UseOp = UseMI.getOpcode();
+ if (UseOp == SPIRV::OpConstantFunctionPointerINTEL ||
+ UseOp == SPIRV::OpEnqueueKernel) {
+ MAI.setSkipEmission(&MI);
+ return false;
+ }
+ }
+ }
for (MachineInstr &UseMI : MRI.use_instructions(DefReg)) {
if (UseMI.getOpcode() != SPIRV::OpConstantFunctionPointerINTEL)
continue;
@@ -563,6 +579,26 @@ void SPIRVModuleAnalysis::collectDeclarations(const Module &M) {
if (DefMO.isReg() && isDeclSection(MRI, MI) &&
!MAI.hasRegisterAlias(MF, DefMO.getReg()))
visitDecl(MRI, SignatureToGReg, GlobalToGReg, MF, MI);
+ // OpEnqueueKernel is not a decl, but its Invoke operand may be a
+ // function-pointer placeholder OpUndef recorded by selectGlobalValue.
+ // Resolve it to the OpFunction's global <id> via visitFunPtrUse.
+ if (Opcode == SPIRV::OpEnqueueKernel && MI.getNumOperands() > 8) {
+ const MachineOperand &InvokeMO = MI.getOperand(8);
+ if (InvokeMO.isReg()) {
+ Register InvokeReg = InvokeMO.getReg();
+ if (!MAI.hasRegisterAlias(MF, InvokeReg)) {
+ if (const MachineInstr *DefMI =
+ MRI.getUniqueVRegDef(InvokeReg)) {
+ if (DefMI->getOpcode() == SPIRV::OpUndef) {
+ const MachineOperand *FunPtrOp = &DefMI->getOperand(0);
+ if (GR->getFunctionDefinitionByUse(FunPtrOp))
+ visitFunPtrUse(InvokeReg, FunPtrOp, SignatureToGReg,
+ GlobalToGReg, MF);
+ }
+ }
+ }
+ }
+ }
}
}
}
@@ -1636,6 +1672,7 @@ void addInstrRequirements(const MachineInstr &MI,
case SPIRV::OpTypeDeviceEvent:
case SPIRV::OpTypeQueue:
case SPIRV::OpBuildNDRange:
+ case SPIRV::OpEnqueueKernel:
Reqs.addCapability(SPIRV::Capability::DeviceEnqueue);
break;
case SPIRV::OpDecorate:
diff --git a/llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll b/llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll
deleted file mode 100644
index 484f86a65880d..0000000000000
--- a/llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll
+++ /dev/null
@@ -1,67 +0,0 @@
-; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-
-; TODO(#60133): Requires updates following opaque pointer migration.
-; XFAIL: *
-
-;; This test checks that Invoke parameter of OpEnueueKernel instruction meet the
-;; following specification requirements in case of enqueueing empty block:
-;; "Invoke must be an OpFunction whose OpTypeFunction operand has:
-;; - Result Type must be OpTypeVoid.
-;; - The first parameter must have a type of OpTypePointer to an 8-bit OpTypeInt.
-;; - An optional list of parameters, each of which must have a type of OpTypePointer to the Workgroup Storage Class.
-;; ... "
-;; __kernel void test_enqueue_empty() {
-;; enqueue_kernel(get_default_queue(),
-;; CLK_ENQUEUE_FLAGS_WAIT_KERNEL,
-;; ndrange_1D(1),
-;; 0, NULL, NULL,
-;; ^(){});
-;; }
-
-%struct.ndrange_t = type { i32, [3 x i64], [3 x i64], [3 x i64] }
-%opencl.queue_t = type opaque
-%opencl.clk_event_t = type opaque
-
- at __block_literal_global = internal addrspace(1) constant { i32, i32 } { i32 8, i32 4 }, align 4
-
-; CHECK-SPIRV: OpName %[[#Block:]] "__block_literal_global"
-; CHECK-SPIRV: %[[#Void:]] = OpTypeVoid
-; CHECK-SPIRV: %[[#Int8:]] = OpTypeInt 8
-; CHECK-SPIRV: %[[#Int8PtrGen:]] = OpTypePointer Generic %[[#Int8]]
-; CHECK-SPIRV: %[[#Int8Ptr:]] = OpTypePointer CrossWorkgroup %[[#Int8]]
-; CHECK-SPIRV: %[[#Block]] = OpVariable %[[#]]
-
-define spir_kernel void @test_enqueue_empty() {
-entry:
- %tmp = alloca %struct.ndrange_t, align 8
- %call = call spir_func ptr @_Z17get_default_queuev()
- call spir_func void @_Z10ndrange_1Dm(ptr sret(ptr) %tmp, i64 1)
- %0 = call i32 @__enqueue_kernel_basic_events(ptr %call, i32 1, ptr %tmp, i32 0, ptr addrspace(4) null, ptr addrspace(4) null, ptr addrspace(4) addrspacecast (ptr @__test_enqueue_empty_block_invoke_kernel to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @__block_literal_global to ptr addrspace(4)))
- ret void
-; CHECK-SPIRV: %[[#Int8PtrBlock:]] = OpBitcast %[[#Int8Ptr]] %[[#Block]]
-; CHECK-SPIRV: %[[#Int8PtrGenBlock:]] = OpPtrCastToGeneric %[[#Int8PtrGen]] %[[#Int8PtrBlock]]
-; CHECK-SPIRV: %[[#]] = OpEnqueueKernel %[[#]] %[[#]] %[[#]] %[[#]] %[[#]] %[[#]] %[[#]] %[[#Invoke:]] %[[#Int8PtrGenBlock]] %[[#]] %[[#]]
-}
-
-declare spir_func ptr @_Z17get_default_queuev()
-
-declare spir_func void @_Z10ndrange_1Dm(ptr sret(ptr), i64)
-
-define internal spir_func void @__test_enqueue_empty_block_invoke(ptr addrspace(4) %.block_descriptor) {
-entry:
- %.block_descriptor.addr = alloca ptr addrspace(4), align 8
- store ptr addrspace(4) %.block_descriptor, ptr %.block_descriptor.addr, align 8
- %block = bitcast ptr addrspace(4) %.block_descriptor to ptr addrspace(4)
- ret void
-}
-
-define internal spir_kernel void @__test_enqueue_empty_block_invoke_kernel(ptr addrspace(4)) {
-entry:
- call void @__test_enqueue_empty_block_invoke(ptr addrspace(4) %0)
- ret void
-}
-
-declare i32 @__enqueue_kernel_basic_events(ptr, i32, ptr, i32, ptr addrspace(4), ptr addrspace(4), ptr addrspace(4), ptr addrspace(4))
-
-; CHECK-SPIRV: %[[#Invoke]] = OpFunction %[[#Void]] None %[[#]]
-; CHECK-SPIRV-NEXT: %[[#]] = OpFunctionParameter %[[#Int8PtrGen]]
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/enqueue_kernel.ll b/llvm/test/CodeGen/SPIRV/transcoding/enqueue_kernel.ll
index 616fa6c5fbafb..2d7e73d71094e 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/enqueue_kernel.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/enqueue_kernel.ll
@@ -1,389 +1,275 @@
-; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
-
-; TODO(#60133): Requires updates following opaque pointer migration.
-; XFAIL: *
-
-; CHECK-SPIRV-DAG: OpEntryPoint Kernel %[[#BlockKer1:]] "__device_side_enqueue_block_invoke_kernel"
-; CHECK-SPIRV-DAG: OpEntryPoint Kernel %[[#BlockKer2:]] "__device_side_enqueue_block_invoke_2_kernel"
-; CHECK-SPIRV-DAG: OpEntryPoint Kernel %[[#BlockKer3:]] "__device_side_enqueue_block_invoke_3_kernel"
-; CHECK-SPIRV-DAG: OpEntryPoint Kernel %[[#BlockKer4:]] "__device_side_enqueue_block_invoke_4_kernel"
-; CHECK-SPIRV-DAG: OpEntryPoint Kernel %[[#BlockKer5:]] "__device_side_enqueue_block_invoke_5_kernel"
-; CHECK-SPIRV-DAG: OpName %[[#BlockGlb1:]] "__block_literal_global"
-; CHECK-SPIRV-DAG: OpName %[[#BlockGlb2:]] "__block_literal_global.1"
-
-; CHECK-SPIRV-DAG: %[[#Int32Ty:]] = OpTypeInt 32
-; CHECK-SPIRV-DAG: %[[#Int8Ty:]] = OpTypeInt 8
-; CHECK-SPIRV-DAG: %[[#VoidTy:]] = OpTypeVoid
-; CHECK-SPIRV-DAG: %[[#Int8PtrGenTy:]] = OpTypePointer Generic %[[#Int8Ty]]
-; CHECK-SPIRV-DAG: %[[#EventTy:]] = OpTypeDeviceEvent
-; CHECK-SPIRV-DAG: %[[#EventPtrTy:]] = OpTypePointer Generic %[[#EventTy]]
-; CHECK-SPIRV-DAG: %[[#Int32LocPtrTy:]] = OpTypePointer Function %[[#Int32Ty]]
-; CHECK-SPIRV-DAG: %[[#BlockStructTy:]] = OpTypeStruct
-; CHECK-SPIRV-DAG: %[[#BlockStructLocPtrTy:]] = OpTypePointer Function %[[#BlockStructTy]]
-; CHECK-SPIRV-DAG: %[[#BlockTy1:]] = OpTypeFunction %[[#VoidTy]] %[[#Int8PtrGenTy]]
-; CHECK-SPIRV-DAG: %[[#BlockTy2:]] = OpTypeFunction %[[#VoidTy]] %[[#Int8PtrGenTy]]
-; CHECK-SPIRV-DAG: %[[#BlockTy3:]] = OpTypeFunction %[[#VoidTy]] %[[#Int8PtrGenTy]]
-
-; CHECK-SPIRV-DAG: %[[#ConstInt0:]] = OpConstantNull %[[#Int32Ty]]
-; CHECK-SPIRV-DAG: %[[#EventNull:]] = OpConstantNull %[[#EventPtrTy]]
-; CHECK-SPIRV-DAG: %[[#ConstInt21:]] = OpConstant %[[#Int32Ty]] 21{{$}}
-; CHECK-SPIRV-DAG: %[[#ConstInt8:]] = OpConstant %[[#Int32Ty]] 8{{$}}
-; CHECK-SPIRV-DAG: %[[#ConstInt24:]] = OpConstant %[[#Int32Ty]] 24{{$}}
-; CHECK-SPIRV-DAG: %[[#ConstInt12:]] = OpConstant %[[#Int32Ty]] 12{{$}}
-; CHECK-SPIRV-DAG: %[[#ConstInt2:]] = OpConstant %[[#Int32Ty]] 2{{$}}
-
-;; typedef struct {int a;} ndrange_t;
-;; #define NULL ((void*)0)
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK
+; RUN: %if spirv-tools %{ llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; CHECK: OpCapability Kernel
+; CHECK-DAG: %[[#typeInt64:]] = OpTypeInt 64 0
+; CHECK-DAG: %[[#typeInt32:]] = OpTypeInt 32 0
+; CHECK-DAG: %[[#typeInt8:]] = OpTypeInt 8 0
+
+; CHECK-DAG: %[[#Num0i32:]] = OpConstantNull %[[#typeInt32]]
+; CHECK-DAG: %[[#Num1i32:]] = OpConstant %[[#typeInt32]] 1
+; CHECK-DAG: %[[#Num2i32:]] = OpConstant %[[#typeInt32]] 2
+; CHECK-DAG: %[[#Num3i32:]] = OpConstant %[[#typeInt32]] 3
+; CHECK-DAG: %[[#Num8i32:]] = OpConstant %[[#typeInt32]] 8
+
+; CHECK-DAG: %[[#Num16i32:]] = OpConstant %[[#typeInt32]] 16
+; CHECK-DAG: %[[#Num29i32:]] = OpConstant %[[#typeInt32]] 29
+; CHECK-DAG: %[[#Num36i32:]] = OpConstant %[[#typeInt32]] 36
+
+; CHECK-DAG: %[[#Num1i64:]] = OpConstant %[[#typeInt64]] 1
+; CHECK-DAG: %[[#Num2i64:]] = OpConstant %[[#typeInt64]] 2
+; CHECK-DAG: %[[#Num4i64:]] = OpConstant %[[#typeInt64]] 4
+
+; CHECK-DAG: %[[#Array3x64:]] = OpTypeArray %[[#typeInt64:]] %[[#Num3i32]]
+; CHECK-DAG: %[[#TypeNDRangeStruct:]] = OpTypeStruct %[[#typeInt32]] %[[#Array3x64]] %[[#Array3x64]] %[[#Array3x64]]
+
+; CHECK-DAG: %[[#pointerInt8:]] = OpTypePointer Generic %[[#typeInt8]]
+; CHECK-DAG: %[[#nullPtrInt8:]] = OpConstantNull %[[#pointerInt8]]
+; CHECK-DAG: %[[#nullArray3x64:]] = OpConstantNull %[[#Array3x64]]
+
+; CHECK-DAG: %[[#typeEvent:]] = OpTypeDeviceEvent
+; CHECK-DAG: %[[#typeEventPtr:]] = OpTypePointer Generic %[[#typeEvent]]
+; CHECK-DAG: %[[#nullPtrEvent:]] = OpConstantNull %[[#typeEventPtr]]
+; CHECK-DAG: %[[#typeVoid:]] = OpTypeVoid
+; CHECK-DAG: %[[#typeWorkgroupPtrInt8:]] = OpTypePointer Workgroup %[[#typeInt8]]
+; CHECK-DAG: %[[#typeFnVoidPtr:]] = OpTypeFunction %[[#typeVoid]] %[[#pointerInt8]]
+; CHECK-DAG: %[[#typeFnVoidPtrLocal1:]] = OpTypeFunction %[[#typeVoid]] %[[#pointerInt8]] %[[#typeWorkgroupPtrInt8]]
+; CHECK-DAG: %[[#typeFnVoidPtrLocal3:]] = OpTypeFunction %[[#typeVoid]] %[[#pointerInt8]] %[[#typeWorkgroupPtrInt8]] %[[#typeWorkgroupPtrInt8]] %[[#typeWorkgroupPtrInt8]]
+
+; CHECK-DAG: OpName %[[#InvokeKernel1:]] "__device_side_enqueue_block_invoke_kernel"
+; CHECK-DAG: OpName %[[#InvokeKernel2:]] "__device_side_enqueue_block_invoke_2_kernel"
+; CHECK-DAG: OpName %[[#InvokeKernel3:]] "__device_side_enqueue_block_invoke_3_kernel"
+; CHECK-DAG: OpName %[[#InvokeKernel4:]] "__device_side_enqueue_block_invoke_4_kernel"
+; CHECK-DAG: OpName %[[#InvokeKernel5:]] "__device_side_enqueue_block_invoke_5_kernel"
+; CHECK-DAG: OpName %[[#InvokeKernel6:]] "__device_side_enqueue_block_invoke_6_kernel"
+
+; CHECK-LABEL: ; -- Begin function device_side_enqueue
+
+; CHECK: %[[#NDRange3sret:]] = OpBuildNDRange %[[#TypeNDRangeStruct]] %[[#]] %[[#]] %[[#]]
+; CHECK-NEXT: OpStore %[[#NDRange3:]] %[[#NDRange3sret]]
+
+;; #define NULL ((void*)0)
;; kernel void device_side_enqueue(global int *a, global int *b, int i, char c0) {
-;; queue_t default_queue;
-;; unsigned flags = 0;
-;; ndrange_t ndrange;
-;; clk_event_t clk_event;
-;; clk_event_t event_wait_list;
-;; clk_event_t event_wait_list2[] = {clk_event};
-
-;; Emits block literal on stack and block kernel.
-
-; CHECK-SPIRV: %[[#BlockLitPtr1:]] = OpBitcast %[[#BlockStructLocPtrTy]]
-; CHECK-SPIRV-NEXT: %[[#BlockLit1:]] = OpPtrCastToGeneric %[[#Int8PtrGenTy]] %[[#BlockLitPtr1]]
-; CHECK-SPIRV-NEXT: %[[#]] = OpEnqueueKernel %[[#Int32Ty]] %[[#]] %[[#]] %[[#]] %[[#ConstInt0]] %[[#EventNull]] %[[#EventNull]] %[[#BlockKer1]] %[[#BlockLit1]] %[[#ConstInt21]] %[[#ConstInt8]]
-
-;; enqueue_kernel(default_queue, flags, ndrange,
-;; ^(void) {
-;; a[i] = c0;
-;; });
-
-;; Emits block literal on stack and block kernel.
-
-; CHECK-SPIRV: %[[#Event1:]] = OpPtrCastToGeneric %[[#EventPtrTy]]
-; CHECK-SPIRV: %[[#Event2:]] = OpPtrCastToGeneric %[[#EventPtrTy]]
-; CHECK-SPIRV: %[[#BlockLitPtr2:]] = OpBitcast %[[#BlockStructLocPtrTy]]
-; CHECK-SPIRV-NEXT: %[[#BlockLit2:]] = OpPtrCastToGeneric %[[#Int8PtrGenTy]] %[[#BlockLitPtr2]]
-; CHECK-SPIRV-NEXT: %[[#]] = OpEnqueueKernel %[[#Int32Ty]] %[[#]] %[[#]] %[[#]] %[[#ConstInt2]] %[[#Event1]] %[[#Event2]] %[[#BlockKer2]] %[[#BlockLit2]] %[[#ConstInt24]] %[[#ConstInt8]]
-
-;; enqueue_kernel(default_queue, flags, ndrange, 2, &event_wait_list, &clk_event,
-;; ^(void) {
-;; a[i] = b[i];
-;; });
-
-;; char c;
-;; Emits global block literal and block kernel.
-
-; CHECK-SPIRV: %[[#Event1:]] = OpPtrCastToGeneric %[[#EventPtrTy]]
-; CHECK-SPIRV: %[[#Event2:]] = OpPtrCastToGeneric %[[#EventPtrTy]]
-; CHECK-SPIRV: %[[#BlockLit3Tmp:]] = OpBitcast %[[#]] %[[#BlockGlb1]]
-; CHECK-SPIRV: %[[#BlockLit3:]] = OpPtrCastToGeneric %[[#Int8PtrGenTy]] %[[#BlockLit3Tmp]]
-; CHECK-SPIRV: %[[#LocalBuf31:]] = OpPtrAccessChain %[[#Int32LocPtrTy]]
-; CHECK-SPIRV: %[[#]] = OpEnqueueKernel %[[#Int32Ty]] %[[#]] %[[#]] %[[#]] %[[#ConstInt2]] %[[#Event1]] %[[#Event2]] %[[#BlockKer3]] %[[#BlockLit3]] %[[#ConstInt12]] %[[#ConstInt8]] %[[#LocalBuf31]]
-
-;; enqueue_kernel(default_queue, flags, ndrange, 2, event_wait_list2, &clk_event,
-;; ^(local void *p) {
-;; return;
-;; },
-;; c);
-
-;; Emits global block literal and block kernel.
-
-; CHECK-SPIRV: %[[#BlockLit4Tmp:]] = OpBitcast %[[#]] %[[#BlockGlb2]]
-; CHECK-SPIRV: %[[#BlockLit4:]] = OpPtrCastToGeneric %[[#Int8PtrGenTy]] %[[#BlockLit4Tmp]]
-; CHECK-SPIRV: %[[#LocalBuf41:]] = OpPtrAccessChain %[[#Int32LocPtrTy]]
-; CHECK-SPIRV-NEXT: %[[#LocalBuf42:]] = OpPtrAccessChain %[[#Int32LocPtrTy]]
-; CHECK-SPIRV-NEXT: %[[#LocalBuf43:]] = OpPtrAccessChain %[[#Int32LocPtrTy]]
-; CHECK-SPIRV-NEXT: %[[#]] = OpEnqueueKernel %[[#Int32Ty]] %[[#]] %[[#]] %[[#]] %[[#ConstInt0]] %[[#EventNull]] %[[#EventNull]] %[[#BlockKer4]] %[[#BlockLit4]] %[[#ConstInt12]] %[[#ConstInt8]] %[[#LocalBuf41]] %[[#LocalBuf42]] %[[#LocalBuf43]]
-
-;; enqueue_kernel(default_queue, flags, ndrange,
-;; ^(local void *p1, local void *p2, local void *p3) {
-;; return;
-;; },
-;; 1, 2, 4);
-
-;; Emits block literal on stack and block kernel.
-
-; CHECK-SPIRV: %[[#Event1:]] = OpPtrCastToGeneric %[[#EventPtrTy]]
-; CHECK-SPIRV: %[[#BlockLit5Tmp:]] = OpBitcast %[[#BlockStructLocPtrTy]]
-; CHECK-SPIRV-NEXT: %[[#BlockLit5:]] = OpPtrCastToGeneric %[[#Int8PtrGenTy]] %[[#BlockLit5Tmp]]
-; CHECK-SPIRV-NEXT: %[[#]] = OpEnqueueKernel %[[#Int32Ty]] %[[#]] %[[#]] %[[#]] %[[#ConstInt0]] %[[#EventNull]] %[[#Event1]] %[[#BlockKer5]] %[[#BlockLit5]] %[[#ConstInt24]] %[[#ConstInt8]]
-
-;; enqueue_kernel(default_queue, flags, ndrange, 0, NULL, &clk_event,
-;; ^(void) {
-;; a[i] = b[i];
-;; });
+;; queue_t default_queue;
+;; unsigned flags = 0;
+;; ndrange_t ndrange;
+;; clk_event_t clk_event;
+;; clk_event_t event_wait_list;
+;; clk_event_t event_wait_list2[] = {clk_event};
+;;
+;; const size_t gs[] = {1,2,4};
+;;
+; TODO: Fix int8 type came from clang in case of NULL events passed in opencl's enqueue_kernel.
+;; // enqueue empty kernel
+; CHECK: %[[#]] = OpEnqueueKernel %[[#typeInt32]] %[[#default_queue:]] %[[#Num1i32]] %[[#NDRange3]] %[[#Num0i32]] %[[#nullPtrInt8]] %[[#nullPtrInt8]] %[[#InvokeKernel1]] %[[#]] %[[#Num16i32]] %[[#Num8i32]]
+;; enqueue_kernel(default_queue,
+;; CLK_ENQUEUE_FLAGS_WAIT_KERNEL,
+;; ndrange_3D(gs),
+;; 0, NULL, NULL,
+;; ^(){});
+;;
+;; // no events, no var args
+; CHECK: %[[#]] = OpEnqueueKernel %[[#typeInt32]] %[[#default_queue]] %[[#Num0i32]] %[[#]] %[[#Num0i32]] %[[#nullPtrEvent]] %[[#nullPtrEvent]] %[[#InvokeKernel2]] %[[#]] %[[#Num29i32]] %[[#Num8i32]]
+;; enqueue_kernel(default_queue, flags, ndrange,
+;; ^(void) {
+;; a[i] = c0;
+;; });
+;;
+;; // event, no var args
+; CHECK: %[[#event1:]] = OpPtrCastToGeneric %[[#typeEventPtr]] %[[#]]
+; CHECK-NEXT: %[[#event2:]] = OpPtrCastToGeneric %[[#typeEventPtr]] %[[#]]
+; CHECK: %[[#]] = OpEnqueueKernel %[[#typeInt32]] %[[#default_queue]] %[[#Num0i32]] %[[#]] %[[#Num2i32]] %[[#event1]] %[[#event2]] %[[#InvokeKernel3]] %[[#]] %[[#Num36i32]] %[[#Num8i32]]
+;; enqueue_kernel(default_queue, flags, ndrange, 2, &event_wait_list, &clk_event,
+;; ^(void) {
+;; a[i] = b[i];
+;; });
+;;
+;; // events, var arg
+; CHECK: %[[#]] = OpEnqueueKernel %[[#typeInt32]] %[[#default_queue]] %[[#Num0i32]] %[[#]] %[[#Num2i32]] %[[#event_wait_list2:]] %[[#event2]] %[[#InvokeKernel4]] %[[#]] %[[#Num16i32]] %[[#Num8i32]] %[[#]]
+;; char c;
+;; enqueue_kernel(default_queue, flags, ndrange, 2, event_wait_list2, &clk_event,
+;; ^(local void *p) {
+;; return;
+;; },
+;; c);
+;;
+;; // no events, three var args
+; CHECK: %[[#]] = OpEnqueueKernel %[[#typeInt32]] %[[#default_queue]] %[[#Num0i32]] %[[#]] %[[#Num0i32]] %[[#nullPtrEvent]] %[[#nullPtrEvent]] %[[#InvokeKernel5]] %[[#]] %[[#Num16i32]] %[[#Num8i32]] %[[#]] %[[#]] %[[#]]
+;; enqueue_kernel(default_queue, flags, ndrange,
+;; ^(local void *p1, local void *p2, local void *p3) {
+;; return;
+;; },
+;; 101, 102, 104);
+;;
+;; // null event, no var args
+; CHECK: %[[#]] = OpEnqueueKernel %[[#typeInt32]] %[[#default_queue]] %[[#Num0i32]] %[[#]] %[[#Num0i32]] %[[#nullPtrInt8]] %[[#event2]] %[[#InvokeKernel6]] %[[#]] %[[#Num36i32]] %[[#Num8i32]]
+;; enqueue_kernel(default_queue, flags, ndrange, 0, NULL, &clk_event,
+;; ^(void) {
+;; a[i] = b[i];
+;; });
;; }
-; CHECK-SPIRV-DAG: %[[#BlockKer1]] = OpFunction %[[#VoidTy]] None %[[#BlockTy1]]
-; CHECK-SPIRV-DAG: %[[#BlockKer2]] = OpFunction %[[#VoidTy]] None %[[#BlockTy1]]
-; CHECK-SPIRV-DAG: %[[#BlockKer3]] = OpFunction %[[#VoidTy]] None %[[#BlockTy3]]
-; CHECK-SPIRV-DAG: %[[#BlockKer4]] = OpFunction %[[#VoidTy]] None %[[#BlockTy2]]
-; CHECK-SPIRV-DAG: %[[#BlockKer5]] = OpFunction %[[#VoidTy]] None %[[#BlockTy1]]
+; CHECK-DAG: %[[#InvokeKernel1]] = OpFunction %[[#typeVoid]] {{Pure|None}} %[[#typeFnVoidPtr]]
+; CHECK-DAG: %[[#InvokeKernel2]] = OpFunction %[[#typeVoid]] {{Pure|None}} %[[#typeFnVoidPtr]]
+; CHECK-DAG: %[[#InvokeKernel3]] = OpFunction %[[#typeVoid]] {{Pure|None}} %[[#typeFnVoidPtr]]
+; CHECK-DAG: %[[#InvokeKernel4]] = OpFunction %[[#typeVoid]] {{Pure|None}} %[[#typeFnVoidPtrLocal1]]
+; CHECK-DAG: %[[#InvokeKernel5]] = OpFunction %[[#typeVoid]] {{Pure|None}} %[[#typeFnVoidPtrLocal3]]
+; CHECK-DAG: %[[#InvokeKernel6]] = OpFunction %[[#typeVoid]] {{Pure|None}} %[[#typeFnVoidPtr]]
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1"
+target triple = "spirv64-unknown-unknown"
-%opencl.queue_t = type opaque
-%struct.ndrange_t = type { i32 }
-%opencl.clk_event_t = type opaque
-%struct.__opencl_block_literal_generic = type { i32, i32, ptr addrspace(4) }
+%struct.ndrange_t = type { i32, [3 x i64], [3 x i64], [3 x i64] }
- at __block_literal_global = internal addrspace(1) constant { i32, i32, ptr addrspace(4) } { i32 12, i32 4, ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_3 to ptr addrspace(4)) }, align 4
- at __block_literal_global.1 = internal addrspace(1) constant { i32, i32, ptr addrspace(4) } { i32 12, i32 4, ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_4 to ptr addrspace(4)) }, align 4
+ at __const.device_side_enqueue.gs = private unnamed_addr addrspace(2) constant [3 x i64] [i64 1, i64 2, i64 4], align 8
+ at __block_literal_global = internal addrspace(1) constant { i32, i32, ptr addrspace(4) } { i32 16, i32 8, ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke to ptr addrspace(4)) }, align 8
+ at __block_literal_global.1 = internal addrspace(1) constant { i32, i32, ptr addrspace(4) } { i32 16, i32 8, ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_4 to ptr addrspace(4)) }, align 8
+ at __block_literal_global.2 = internal addrspace(1) constant { i32, i32, ptr addrspace(4) } { i32 16, i32 8, ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_5 to ptr addrspace(4)) }, align 8
-define dso_local spir_kernel void @device_side_enqueue(ptr addrspace(1) noundef %a, ptr addrspace(1) noundef %b, i32 noundef %i, i8 noundef signext %c0) {
+define spir_kernel void @device_side_enqueue(ptr addrspace(1) align 4 %a, ptr addrspace(1) align 4 %b, i32 %i, i8 %c0, target("spirv.Queue") %default_queue) {
entry:
- %a.addr = alloca ptr addrspace(1), align 4
- %b.addr = alloca ptr addrspace(1), align 4
- %i.addr = alloca i32, align 4
- %c0.addr = alloca i8, align 1
- %default_queue = alloca ptr, align 4
- %flags = alloca i32, align 4
- %ndrange = alloca %struct.ndrange_t, align 4
- %clk_event = alloca ptr, align 4
- %event_wait_list = alloca ptr, align 4
- %event_wait_list2 = alloca [1 x ptr], align 4
- %tmp = alloca %struct.ndrange_t, align 4
- %block = alloca <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, i8 }>, align 4
- %tmp3 = alloca %struct.ndrange_t, align 4
- %block4 = alloca <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, align 4
- %c = alloca i8, align 1
- %tmp11 = alloca %struct.ndrange_t, align 4
- %block_sizes = alloca [1 x i32], align 4
- %tmp12 = alloca %struct.ndrange_t, align 4
- %block_sizes13 = alloca [3 x i32], align 4
- %tmp14 = alloca %struct.ndrange_t, align 4
- %block15 = alloca <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, align 4
- store ptr addrspace(1) %a, ptr %a.addr, align 4
- store ptr addrspace(1) %b, ptr %b.addr, align 4
- store i32 %i, ptr %i.addr, align 4
- store i8 %c0, ptr %c0.addr, align 1
- store i32 0, ptr %flags, align 4
- %arrayinit.begin = getelementptr inbounds [1 x ptr], ptr %event_wait_list2, i32 0, i32 0
- %0 = load ptr, ptr %clk_event, align 4
- store ptr %0, ptr %arrayinit.begin, align 4
- %1 = load ptr, ptr %default_queue, align 4
- %2 = load i32, ptr %flags, align 4
- %3 = bitcast ptr %tmp to ptr
- %4 = bitcast ptr %ndrange to ptr
- call void @llvm.memcpy.p0.p0.i32(ptr align 4 %3, ptr align 4 %4, i32 4, i1 false)
- %block.size = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, i8 }>, ptr %block, i32 0, i32 0
- store i32 21, ptr %block.size, align 4
- %block.align = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, i8 }>, ptr %block, i32 0, i32 1
- store i32 4, ptr %block.align, align 4
- %block.invoke = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, i8 }>, ptr %block, i32 0, i32 2
- store ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke to ptr addrspace(4)), ptr %block.invoke, align 4
- %block.captured = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, i8 }>, ptr %block, i32 0, i32 3
- %5 = load ptr addrspace(1), ptr %a.addr, align 4
- store ptr addrspace(1) %5, ptr %block.captured, align 4
- %block.captured1 = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, i8 }>, ptr %block, i32 0, i32 4
- %6 = load i32, ptr %i.addr, align 4
- store i32 %6, ptr %block.captured1, align 4
- %block.captured2 = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, i8 }>, ptr %block, i32 0, i32 5
- %7 = load i8, ptr %c0.addr, align 1
- store i8 %7, ptr %block.captured2, align 4
- %8 = bitcast ptr %block to ptr
- %9 = addrspacecast ptr %8 to ptr addrspace(4)
- %10 = call spir_func i32 @__enqueue_kernel_basic(ptr %1, i32 %2, ptr byval(%struct.ndrange_t) %tmp, ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_kernel to ptr addrspace(4)), ptr addrspace(4) %9)
- %11 = load ptr, ptr %default_queue, align 4
- %12 = load i32, ptr %flags, align 4
- %13 = bitcast ptr %tmp3 to ptr
- %14 = bitcast ptr %ndrange to ptr
- call void @llvm.memcpy.p0.p0.i32(ptr align 4 %13, ptr align 4 %14, i32 4, i1 false)
- %15 = addrspacecast ptr %event_wait_list to ptr addrspace(4)
- %16 = addrspacecast ptr %clk_event to ptr addrspace(4)
- %block.size5 = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr %block4, i32 0, i32 0
- store i32 24, ptr %block.size5, align 4
- %block.align6 = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr %block4, i32 0, i32 1
- store i32 4, ptr %block.align6, align 4
- %block.invoke7 = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr %block4, i32 0, i32 2
- store ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_2 to ptr addrspace(4)), ptr %block.invoke7, align 4
- %block.captured8 = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr %block4, i32 0, i32 3
- %17 = load ptr addrspace(1), ptr %a.addr, align 4
- store ptr addrspace(1) %17, ptr %block.captured8, align 4
- %block.captured9 = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr %block4, i32 0, i32 4
- %18 = load i32, ptr %i.addr, align 4
- store i32 %18, ptr %block.captured9, align 4
- %block.captured10 = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr %block4, i32 0, i32 5
- %19 = load ptr addrspace(1), ptr %b.addr, align 4
- store ptr addrspace(1) %19, ptr %block.captured10, align 4
- %20 = bitcast ptr %block4 to ptr
- %21 = addrspacecast ptr %20 to ptr addrspace(4)
- %22 = call spir_func i32 @__enqueue_kernel_basic_events(ptr %11, i32 %12, ptr %tmp3, i32 2, ptr addrspace(4) %15, ptr addrspace(4) %16, ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_2_kernel to ptr addrspace(4)), ptr addrspace(4) %21)
- %23 = load ptr, ptr %default_queue, align 4
- %24 = load i32, ptr %flags, align 4
- %25 = bitcast ptr %tmp11 to ptr
- %26 = bitcast ptr %ndrange to ptr
- call void @llvm.memcpy.p0.p0.i32(ptr align 4 %25, ptr align 4 %26, i32 4, i1 false)
- %arraydecay = getelementptr inbounds [1 x ptr], ptr %event_wait_list2, i32 0, i32 0
- %27 = addrspacecast ptr %arraydecay to ptr addrspace(4)
- %28 = addrspacecast ptr %clk_event to ptr addrspace(4)
- %29 = getelementptr [1 x i32], ptr %block_sizes, i32 0, i32 0
- %30 = load i8, ptr %c, align 1
- %31 = zext i8 %30 to i32
- store i32 %31, ptr %29, align 4
- %32 = call spir_func i32 @__enqueue_kernel_events_varargs(ptr %23, i32 %24, ptr %tmp11, i32 2, ptr addrspace(4) %27, ptr addrspace(4) %28, ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_3_kernel to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @__block_literal_global to ptr addrspace(4)), i32 1, ptr %29)
- %33 = load ptr, ptr %default_queue, align 4
- %34 = load i32, ptr %flags, align 4
- %35 = bitcast ptr %tmp12 to ptr
- %36 = bitcast ptr %ndrange to ptr
- call void @llvm.memcpy.p0.p0.i32(ptr align 4 %35, ptr align 4 %36, i32 4, i1 false)
- %37 = getelementptr [3 x i32], ptr %block_sizes13, i32 0, i32 0
- store i32 1, ptr %37, align 4
- %38 = getelementptr [3 x i32], ptr %block_sizes13, i32 0, i32 1
- store i32 2, ptr %38, align 4
- %39 = getelementptr [3 x i32], ptr %block_sizes13, i32 0, i32 2
- store i32 4, ptr %39, align 4
- %40 = call spir_func i32 @__enqueue_kernel_varargs(ptr %33, i32 %34, ptr %tmp12, ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_4_kernel to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @__block_literal_global.1 to ptr addrspace(4)), i32 3, ptr %37)
- %41 = load ptr, ptr %default_queue, align 4
- %42 = load i32, ptr %flags, align 4
- %43 = bitcast ptr %tmp14 to ptr
- %44 = bitcast ptr %ndrange to ptr
- call void @llvm.memcpy.p0.p0.i32(ptr align 4 %43, ptr align 4 %44, i32 4, i1 false)
- %45 = addrspacecast ptr %clk_event to ptr addrspace(4)
- %block.size16 = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr %block15, i32 0, i32 0
- store i32 24, ptr %block.size16, align 4
- %block.align17 = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr %block15, i32 0, i32 1
- store i32 4, ptr %block.align17, align 4
- %block.invoke18 = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr %block15, i32 0, i32 2
- store ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_5 to ptr addrspace(4)), ptr %block.invoke18, align 4
- %block.captured19 = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr %block15, i32 0, i32 3
- %46 = load ptr addrspace(1), ptr %a.addr, align 4
- store ptr addrspace(1) %46, ptr %block.captured19, align 4
- %block.captured20 = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr %block15, i32 0, i32 4
- %47 = load i32, ptr %i.addr, align 4
- store i32 %47, ptr %block.captured20, align 4
- %block.captured21 = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr %block15, i32 0, i32 5
- %48 = load ptr addrspace(1), ptr %b.addr, align 4
- store ptr addrspace(1) %48, ptr %block.captured21, align 4
- %49 = bitcast ptr %block15 to ptr
- %50 = addrspacecast ptr %49 to ptr addrspace(4)
- %51 = call spir_func i32 @__enqueue_kernel_basic_events(ptr %41, i32 %42, ptr %tmp14, i32 0, ptr addrspace(4) null, ptr addrspace(4) %45, ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_5_kernel to ptr addrspace(4)), ptr addrspace(4) %50)
+ %clk_event.i = alloca target("spirv.DeviceEvent"), align 8
+ %event_wait_list.i = alloca target("spirv.DeviceEvent"), align 8
+ %event_wait_list2.i = alloca [1 x target("spirv.DeviceEvent")], align 8
+ %gs.i = alloca [3 x i64], align 8
+ %tmp.i = alloca %struct.ndrange_t, align 8
+ %tmp1.i = alloca %struct.ndrange_t, align 8
+ %block.i = alloca <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, i8 }>, align 8
+ %tmp4.i = alloca %struct.ndrange_t, align 8
+ %block5.i = alloca <{ i32, i32, ptr addrspace(4), ptr addrspace(1), ptr addrspace(1), i32 }>, align 8
+ %tmp12.i = alloca %struct.ndrange_t, align 8
+ %block_sizes.i = alloca [1 x i64], align 8
+ %tmp14.i = alloca %struct.ndrange_t, align 8
+ %block_sizes15.i = alloca [3 x i64], align 8
+ %tmp16.i = alloca %struct.ndrange_t, align 8
+ %block17.i = alloca <{ i32, i32, ptr addrspace(4), ptr addrspace(1), ptr addrspace(1), i32 }>, align 8
+ call void @llvm.memcpy.p0.p2.i64(ptr align 8 %gs.i, ptr addrspace(2) align 8 @__const.device_side_enqueue.gs, i64 24, i1 false)
+ call spir_func void @_Z10ndrange_3DPKm(ptr sret(%struct.ndrange_t) align 8 %tmp.i, ptr %gs.i)
+ %0 = call spir_func i32 @__enqueue_kernel_basic_events(target("spirv.Queue") %default_queue, i32 1, ptr %tmp.i, i32 0, ptr addrspace(4) null, ptr addrspace(4) null, ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_kernel to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @__block_literal_global to ptr addrspace(4)))
+ store i32 29, ptr %block.i, align 8
+ %block.align.i = getelementptr inbounds i8, ptr %block.i, i64 4
+ store i32 8, ptr %block.align.i, align 4
+ %block.invoke.i = getelementptr inbounds i8, ptr %block.i, i64 8
+ store ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_2 to ptr addrspace(4)), ptr %block.invoke.i, align 8
+ %block.captured.i = getelementptr inbounds i8, ptr %block.i, i64 16
+ store ptr addrspace(1) %a, ptr %block.captured.i, align 8
+ %block.captured2.i = getelementptr inbounds i8, ptr %block.i, i64 24
+ store i32 %i, ptr %block.captured2.i, align 8
+ %block.captured3.i = getelementptr inbounds i8, ptr %block.i, i64 28
+ store i8 %c0, ptr %block.captured3.i, align 4
+ %1 = addrspacecast ptr %block.i to ptr addrspace(4)
+ %2 = call spir_func i32 @__enqueue_kernel_basic(target("spirv.Queue") %default_queue, i32 0, ptr %tmp1.i, ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_2_kernel to ptr addrspace(4)), ptr addrspace(4) %1)
+ %3 = addrspacecast ptr %event_wait_list.i to ptr addrspace(4)
+ %4 = addrspacecast ptr %clk_event.i to ptr addrspace(4)
+ store i32 36, ptr %block5.i, align 8
+ %block.align7.i = getelementptr inbounds i8, ptr %block5.i, i64 4
+ store i32 8, ptr %block.align7.i, align 4
+ %block.invoke8.i = getelementptr inbounds i8, ptr %block5.i, i64 8
+ store ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_3 to ptr addrspace(4)), ptr %block.invoke8.i, align 8
+ %block.captured9.i = getelementptr inbounds i8, ptr %block5.i, i64 16
+ store ptr addrspace(1) %a, ptr %block.captured9.i, align 8
+ %block.captured10.i = getelementptr inbounds i8, ptr %block5.i, i64 32
+ store i32 %i, ptr %block.captured10.i, align 8
+ %block.captured11.i = getelementptr inbounds i8, ptr %block5.i, i64 24
+ store ptr addrspace(1) %b, ptr %block.captured11.i, align 8
+ %5 = addrspacecast ptr %block5.i to ptr addrspace(4)
+ %6 = call spir_func i32 @__enqueue_kernel_basic_events(target("spirv.Queue") %default_queue, i32 0, ptr %tmp4.i, i32 2, ptr addrspace(4) %3, ptr addrspace(4) %4, ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_3_kernel to ptr addrspace(4)), ptr addrspace(4) %5)
+ %7 = addrspacecast ptr %event_wait_list2.i to ptr addrspace(4)
+ store i64 0, ptr %block_sizes.i, align 8
+ %8 = call spir_func i32 @__enqueue_kernel_events_varargs(target("spirv.Queue") %default_queue, i32 0, ptr %tmp12.i, i32 2, ptr addrspace(4) %7, ptr addrspace(4) %4, ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_4_kernel to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @__block_literal_global.1 to ptr addrspace(4)), i32 1, ptr %block_sizes.i)
+ store i64 101, ptr %block_sizes15.i, align 8
+ %9 = getelementptr inbounds i8, ptr %block_sizes15.i, i64 8
+ store i64 102, ptr %9, align 8
+ %10 = getelementptr inbounds i8, ptr %block_sizes15.i, i64 16
+ store i64 104, ptr %10, align 8
+ %11 = call spir_func i32 @__enqueue_kernel_varargs(target("spirv.Queue") %default_queue, i32 0, ptr %tmp14.i, ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_5_kernel to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @__block_literal_global.2 to ptr addrspace(4)), i32 3, ptr %block_sizes15.i)
+ store i32 36, ptr %block17.i, align 8
+ %block.align19.i = getelementptr inbounds i8, ptr %block17.i, i64 4
+ store i32 8, ptr %block.align19.i, align 4
+ %block.invoke20.i = getelementptr inbounds i8, ptr %block17.i, i64 8
+ store ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_6 to ptr addrspace(4)), ptr %block.invoke20.i, align 8
+ %block.captured21.i = getelementptr inbounds i8, ptr %block17.i, i64 16
+ store ptr addrspace(1) %a, ptr %block.captured21.i, align 8
+ %block.captured22.i = getelementptr inbounds i8, ptr %block17.i, i64 32
+ store i32 %i, ptr %block.captured22.i, align 8
+ %block.captured23.i = getelementptr inbounds i8, ptr %block17.i, i64 24
+ store ptr addrspace(1) %b, ptr %block.captured23.i, align 8
+ %12 = addrspacecast ptr %block17.i to ptr addrspace(4)
+ %13 = call spir_func i32 @__enqueue_kernel_basic_events(target("spirv.Queue") %default_queue, i32 0, ptr %tmp16.i, i32 0, ptr addrspace(4) null, ptr addrspace(4) %4, ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_6_kernel to ptr addrspace(4)), ptr addrspace(4) %12)
ret void
}
-declare void @llvm.memcpy.p0.p0.i32(ptr noalias nocapture writeonly, ptr noalias nocapture readonly, i32, i1 immarg)
+declare void @llvm.lifetime.start.p0(ptr)
+
+declare void @llvm.memcpy.p0.p2.i64(ptr, ptr addrspace(2), i64, i1 immarg)
+
+declare spir_func void @_Z10ndrange_3DPKm(ptr sret(%struct.ndrange_t) align 8, ptr)
-define internal spir_func void @__device_side_enqueue_block_invoke(ptr addrspace(4) noundef %.block_descriptor) {
+define internal spir_func void @__device_side_enqueue_block_invoke(ptr addrspace(4) %.block_descriptor) {
entry:
- %.block_descriptor.addr = alloca ptr addrspace(4), align 4
- %block.addr = alloca ptr addrspace(4), align 4
- store ptr addrspace(4) %.block_descriptor, ptr %.block_descriptor.addr, align 4
- %block = bitcast ptr addrspace(4) %.block_descriptor to ptr addrspace(4)
- store ptr addrspace(4) %block, ptr %block.addr, align 4
- %block.capture.addr = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, i8 }>, ptr addrspace(4) %block, i32 0, i32 5
- %0 = load i8, ptr addrspace(4) %block.capture.addr, align 4
- %conv = sext i8 %0 to i32
- %block.capture.addr1 = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, i8 }>, ptr addrspace(4) %block, i32 0, i32 3
- %1 = load ptr addrspace(1), ptr addrspace(4) %block.capture.addr1, align 4
- %block.capture.addr2 = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, i8 }>, ptr addrspace(4) %block, i32 0, i32 4
- %2 = load i32, ptr addrspace(4) %block.capture.addr2, align 4
- %arrayidx = getelementptr inbounds i32, ptr addrspace(1) %1, i32 %2
- store i32 %conv, ptr addrspace(1) %arrayidx, align 4
ret void
}
-define spir_kernel void @__device_side_enqueue_block_invoke_kernel(ptr addrspace(4) %0) {
+define internal spir_kernel void @__device_side_enqueue_block_invoke_kernel(ptr addrspace(4) %0) {
entry:
- call spir_func void @__device_side_enqueue_block_invoke(ptr addrspace(4) %0)
ret void
}
-declare spir_func i32 @__enqueue_kernel_basic(ptr, i32, ptr, ptr addrspace(4), ptr addrspace(4))
+declare spir_func i32 @__enqueue_kernel_basic_events(target("spirv.Queue"), i32, ptr, i32, ptr addrspace(4), ptr addrspace(4), ptr addrspace(4), ptr addrspace(4))
-define internal spir_func void @__device_side_enqueue_block_invoke_2(ptr addrspace(4) noundef %.block_descriptor) {
-entry:
- %.block_descriptor.addr = alloca ptr addrspace(4), align 4
- %block.addr = alloca ptr addrspace(4), align 4
- store ptr addrspace(4) %.block_descriptor, ptr %.block_descriptor.addr, align 4
- %block = bitcast ptr addrspace(4) %.block_descriptor to ptr addrspace(4)
- store ptr addrspace(4) %block, ptr %block.addr, align 4
- %block.capture.addr = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr addrspace(4) %block, i32 0, i32 5
- %0 = load ptr addrspace(1), ptr addrspace(4) %block.capture.addr, align 4
- %block.capture.addr1 = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr addrspace(4) %block, i32 0, i32 4
- %1 = load i32, ptr addrspace(4) %block.capture.addr1, align 4
- %arrayidx = getelementptr inbounds i32, ptr addrspace(1) %0, i32 %1
- %2 = load i32, ptr addrspace(1) %arrayidx, align 4
- %block.capture.addr2 = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr addrspace(4) %block, i32 0, i32 3
- %3 = load ptr addrspace(1), ptr addrspace(4) %block.capture.addr2, align 4
- %block.capture.addr3 = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr addrspace(4) %block, i32 0, i32 4
- %4 = load i32, ptr addrspace(4) %block.capture.addr3, align 4
- %arrayidx4 = getelementptr inbounds i32, ptr addrspace(1) %3, i32 %4
- store i32 %2, ptr addrspace(1) %arrayidx4, align 4
+define internal spir_func void @__device_side_enqueue_block_invoke_2(ptr addrspace(4) %.block_descriptor) {
ret void
}
-define spir_kernel void @__device_side_enqueue_block_invoke_2_kernel(ptr addrspace(4) %0) {
-entry:
- call spir_func void @__device_side_enqueue_block_invoke_2(ptr addrspace(4) %0)
+define internal spir_kernel void @__device_side_enqueue_block_invoke_2_kernel(ptr addrspace(4) %0) {
ret void
}
-declare spir_func i32 @__enqueue_kernel_basic_events(ptr, i32, ptr, i32, ptr addrspace(4), ptr addrspace(4), ptr addrspace(4), ptr addrspace(4))
+declare spir_func i32 @__enqueue_kernel_basic(target("spirv.Queue"), i32, ptr, ptr addrspace(4), ptr addrspace(4))
+
+define internal spir_func void @__device_side_enqueue_block_invoke_3(ptr addrspace(4) %.block_descriptor) {
+ ret void
+}
+
+define internal spir_kernel void @__device_side_enqueue_block_invoke_3_kernel(ptr addrspace(4) %0) {
+ ret void
+}
-define internal spir_func void @__device_side_enqueue_block_invoke_3(ptr addrspace(4) noundef %.block_descriptor, ptr addrspace(3) noundef %p) {
+define internal spir_func void @__device_side_enqueue_block_invoke_4(ptr addrspace(4) %.block_descriptor, ptr addrspace(3) %p) {
entry:
- %.block_descriptor.addr = alloca ptr addrspace(4), align 4
- %p.addr = alloca ptr addrspace(3), align 4
- %block.addr = alloca ptr addrspace(4), align 4
- store ptr addrspace(4) %.block_descriptor, ptr %.block_descriptor.addr, align 4
- %block = bitcast ptr addrspace(4) %.block_descriptor to ptr addrspace(4)
- store ptr addrspace(3) %p, ptr %p.addr, align 4
- store ptr addrspace(4) %block, ptr %block.addr, align 4
ret void
}
-define spir_kernel void @__device_side_enqueue_block_invoke_3_kernel(ptr addrspace(4) %0, ptr addrspace(3) %1) {
+define internal spir_kernel void @__device_side_enqueue_block_invoke_4_kernel(ptr addrspace(4) %0, ptr addrspace(3) %1) {
entry:
- call spir_func void @__device_side_enqueue_block_invoke_3(ptr addrspace(4) %0, ptr addrspace(3) %1)
ret void
}
-declare spir_func i32 @__enqueue_kernel_events_varargs(ptr, i32, ptr, i32, ptr addrspace(4), ptr addrspace(4), ptr addrspace(4), ptr addrspace(4), i32, ptr)
+declare spir_func i32 @__enqueue_kernel_events_varargs(target("spirv.Queue"), i32, ptr, i32, ptr addrspace(4), ptr addrspace(4), ptr addrspace(4), ptr addrspace(4), i32, ptr)
+
+declare void @llvm.lifetime.end.p0(ptr)
-define internal spir_func void @__device_side_enqueue_block_invoke_4(ptr addrspace(4) noundef %.block_descriptor, ptr addrspace(3) noundef %p1, ptr addrspace(3) noundef %p2, ptr addrspace(3) noundef %p3) {
+define internal spir_func void @__device_side_enqueue_block_invoke_5(ptr addrspace(4) %.block_descriptor, ptr addrspace(3) %p1, ptr addrspace(3) %p2, ptr addrspace(3) %p3) {
entry:
- %.block_descriptor.addr = alloca ptr addrspace(4), align 4
- %p1.addr = alloca ptr addrspace(3), align 4
- %p2.addr = alloca ptr addrspace(3), align 4
- %p3.addr = alloca ptr addrspace(3), align 4
- %block.addr = alloca ptr addrspace(4), align 4
- store ptr addrspace(4) %.block_descriptor, ptr %.block_descriptor.addr, align 4
- %block = bitcast ptr addrspace(4) %.block_descriptor to ptr addrspace(4)
- store ptr addrspace(3) %p1, ptr %p1.addr, align 4
- store ptr addrspace(3) %p2, ptr %p2.addr, align 4
- store ptr addrspace(3) %p3, ptr %p3.addr, align 4
- store ptr addrspace(4) %block, ptr %block.addr, align 4
ret void
}
-define spir_kernel void @__device_side_enqueue_block_invoke_4_kernel(ptr addrspace(4) %0, ptr addrspace(3) %1, ptr addrspace(3) %2, ptr addrspace(3) %3) {
+define internal spir_kernel void @__device_side_enqueue_block_invoke_5_kernel(ptr addrspace(4) %0, ptr addrspace(3) %1, ptr addrspace(3) %2, ptr addrspace(3) %3) {
entry:
- call spir_func void @__device_side_enqueue_block_invoke_4(ptr addrspace(4) %0, ptr addrspace(3) %1, ptr addrspace(3) %2, ptr addrspace(3) %3)
ret void
}
-declare spir_func i32 @__enqueue_kernel_varargs(ptr, i32, ptr, ptr addrspace(4), ptr addrspace(4), i32, ptr)
+declare spir_func i32 @__enqueue_kernel_varargs(target("spirv.Queue"), i32, ptr, ptr addrspace(4), ptr addrspace(4), i32, ptr)
-define internal spir_func void @__device_side_enqueue_block_invoke_5(ptr addrspace(4) noundef %.block_descriptor) {
-entry:
- %.block_descriptor.addr = alloca ptr addrspace(4), align 4
- %block.addr = alloca ptr addrspace(4), align 4
- store ptr addrspace(4) %.block_descriptor, ptr %.block_descriptor.addr, align 4
- %block = bitcast ptr addrspace(4) %.block_descriptor to ptr addrspace(4)
- store ptr addrspace(4) %block, ptr %block.addr, align 4
- %block.capture.addr = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr addrspace(4) %block, i32 0, i32 5
- %0 = load ptr addrspace(1), ptr addrspace(4) %block.capture.addr, align 4
- %block.capture.addr1 = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr addrspace(4) %block, i32 0, i32 4
- %1 = load i32, ptr addrspace(4) %block.capture.addr1, align 4
- %arrayidx = getelementptr inbounds i32, ptr addrspace(1) %0, i32 %1
- %2 = load i32, ptr addrspace(1) %arrayidx, align 4
- %block.capture.addr2 = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr addrspace(4) %block, i32 0, i32 3
- %3 = load ptr addrspace(1), ptr addrspace(4) %block.capture.addr2, align 4
- %block.capture.addr3 = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr addrspace(4) %block, i32 0, i32 4
- %4 = load i32, ptr addrspace(4) %block.capture.addr3, align 4
- %arrayidx4 = getelementptr inbounds i32, ptr addrspace(1) %3, i32 %4
- store i32 %2, ptr addrspace(1) %arrayidx4, align 4
+define internal spir_func void @__device_side_enqueue_block_invoke_6(ptr addrspace(4) %.block_descriptor) {
ret void
}
-define spir_kernel void @__device_side_enqueue_block_invoke_5_kernel(ptr addrspace(4) %0) {
-entry:
- call spir_func void @__device_side_enqueue_block_invoke_5(ptr addrspace(4) %0)
+define internal spir_kernel void @__device_side_enqueue_block_invoke_6_kernel(ptr addrspace(4) %0) {
ret void
}
+
+
+!opencl.ocl.version = !{!0}
+
+!0 = !{i32 3, i32 0}
More information about the llvm-commits
mailing list