[llvm] [SPIRV] Fix enqueue empty kernel (PR #187671)
via llvm-commits
llvm-commits at lists.llvm.org
Tue Apr 7 11:14:11 PDT 2026
https://github.com/idubinov updated https://github.com/llvm/llvm-project/pull/187671
>From 69819fc2e6da474b24285af798434fe615430edd Mon Sep 17 00:00:00 2001
From: idubinov <igor.dubinov at amd.com>
Date: Fri, 20 Mar 2026 05:22:22 -0500
Subject: [PATCH 01/12] fix enqueue empty kernel
---
llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp | 52 ++++++++++++++-----
llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll | 7 ++-
2 files changed, 43 insertions(+), 16 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
index f9a9127446013..5020b6aa30a63 100644
--- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
@@ -375,14 +375,24 @@ static MachineInstr *getBlockStructInstr(Register ParamReg,
// 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)
+ //
+ // For function pointers (empty kernel blocks), the sequence may be simpler:
+ // %0:_(pN) = G_GLOBAL_VALUE @block_invoke_kernel
+ // %1:_(p4) = G_ADDRSPACE_CAST %0:_(pN)
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();
+ Register CastSourceReg = MI->getOperand(1).getReg();
+ MachineInstr *SourceMI = MRI->getUniqueVRegDef(CastSourceReg);
+
+ // Check if it's a direct G_GLOBAL_VALUE (function pointer case)
+ if (SourceMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
+ return SourceMI;
+
+ // Otherwise, expect the bitcast sequence (block literal case)
+ assert(isSpvIntrinsic(*SourceMI, Intrinsic::spv_bitcast) &&
+ SourceMI->getOperand(2).isReg());
+ Register ValueReg = SourceMI->getOperand(2).getReg();
MachineInstr *ValueMI = MRI->getUniqueVRegDef(ValueReg);
return ValueMI;
}
@@ -2737,13 +2747,35 @@ static bool buildEnqueueKernel(const SPIRV::IncomingCall *Call,
}
}
+ // Prepare block invoke function and block literal before building OpEnqueueKernel.
+ const unsigned BlockFIdx = HasEvents ? 6 : 3;
+ MachineInstr *BlockMI = getBlockStructInstr(Call->Arguments[BlockFIdx], MRI);
+ assert(BlockMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE);
+
+ Register BlockLiteralReg = Call->Arguments[BlockFIdx + 1];
+ Type *PType = const_cast<Type *>(getBlockStructType(BlockLiteralReg, MRI));
+
+ // OpEnqueueKernel requires the Param to be a pointer to i8 (per SPIR-V spec).
+ // BlockLiteralReg is a Generic pointer to the block struct.
+ // Bitcast it to a Generic pointer to i8.
+ const SPIRVTypeInst Int8Ty = GR->getOrCreateSPIRVIntegerType(8, MIRBuilder);
+ const SPIRVTypeInst Int8PtrGen =
+ GR->getOrCreateSPIRVPointerType(Int8Ty, MIRBuilder,
+ SPIRV::StorageClass::Generic);
+
+ Register BlockLiteralGenAsI8 =
+ createVirtualRegister(Int8PtrGen, GR, MIRBuilder);
+ MIRBuilder.buildInstr(SPIRV::OpBitcast)
+ .addDef(BlockLiteralGenAsI8)
+ .addUse(GR->getSPIRVTypeID(Int8PtrGen))
+ .addUse(BlockLiteralReg);
+
// SPIRV OpEnqueueKernel instruction has 10+ 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]);
@@ -2756,16 +2788,12 @@ static bool buildEnqueueKernel(const SPIRV::IncomingCall *Call,
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);
+ // Param: Pointer to block literal (as Generic i8*).
+ MIB.addUse(BlockLiteralGenAsI8);
- 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));
diff --git a/llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll b/llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll
index 484f86a65880d..442f0dc3cf486 100644
--- a/llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll
+++ b/llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll
@@ -1,7 +1,7 @@
; 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:
@@ -28,7 +28,6 @@
; 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() {
@@ -38,8 +37,8 @@ entry:
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: %[[#BlockGeneric:]] = OpPtrCastToGeneric %[[#]] %[[#Block]]
+; CHECK-SPIRV: %[[#Int8PtrGenBlock:]] = OpBitcast %[[#Int8PtrGen]] %[[#BlockGeneric]]
; CHECK-SPIRV: %[[#]] = OpEnqueueKernel %[[#]] %[[#]] %[[#]] %[[#]] %[[#]] %[[#]] %[[#]] %[[#Invoke:]] %[[#Int8PtrGenBlock]] %[[#]] %[[#]]
}
>From 0ff1cfe00220215d2596ef5d2d22c9b4cccdfac9 Mon Sep 17 00:00:00 2001
From: idubinov <igor.dubinov at amd.com>
Date: Fri, 20 Mar 2026 05:29:30 -0500
Subject: [PATCH 02/12] code format
---
llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp | 8 ++++----
1 file changed, 4 insertions(+), 4 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
index 5020b6aa30a63..bb88de34aa7f0 100644
--- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
@@ -2747,7 +2747,8 @@ static bool buildEnqueueKernel(const SPIRV::IncomingCall *Call,
}
}
- // Prepare block invoke function and block literal before building OpEnqueueKernel.
+ // Prepare block invoke function and block literal before building
+ // OpEnqueueKernel.
const unsigned BlockFIdx = HasEvents ? 6 : 3;
MachineInstr *BlockMI = getBlockStructInstr(Call->Arguments[BlockFIdx], MRI);
assert(BlockMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE);
@@ -2759,9 +2760,8 @@ static bool buildEnqueueKernel(const SPIRV::IncomingCall *Call,
// BlockLiteralReg is a Generic pointer to the block struct.
// Bitcast it to a Generic pointer to i8.
const SPIRVTypeInst Int8Ty = GR->getOrCreateSPIRVIntegerType(8, MIRBuilder);
- const SPIRVTypeInst Int8PtrGen =
- GR->getOrCreateSPIRVPointerType(Int8Ty, MIRBuilder,
- SPIRV::StorageClass::Generic);
+ const SPIRVTypeInst Int8PtrGen = GR->getOrCreateSPIRVPointerType(
+ Int8Ty, MIRBuilder, SPIRV::StorageClass::Generic);
Register BlockLiteralGenAsI8 =
createVirtualRegister(Int8PtrGen, GR, MIRBuilder);
>From 1faa1e0a54c8a90f7dd933ebf9ecce9fb6e797a2 Mon Sep 17 00:00:00 2001
From: idubinov <igor.dubinov at amd.com>
Date: Fri, 20 Mar 2026 06:22:02 -0500
Subject: [PATCH 03/12] fix enqueue kernel test
---
llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp | 49 ++++++++++++++----
llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp | 1 +
.../SPIRV/transcoding/enqueue_kernel.ll | 50 +++++--------------
3 files changed, 52 insertions(+), 48 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
index bb88de34aa7f0..101222651e330 100644
--- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
@@ -370,13 +370,23 @@ lookupBuiltin(StringRef DemangledCall,
static MachineInstr *getBlockStructInstr(Register ParamReg,
MachineRegisterInfo *MRI) {
- // We expect the following sequence of instructions:
+ // We expect one of the following sequences of instructions:
+ //
+ // 1. Stack-allocated blocks (with bitcast):
// %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)
//
- // For function pointers (empty kernel blocks), the sequence may be simpler:
+ // 2. Stack-allocated blocks (direct):
+ // %0:_(pN) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.alloca)
+ // %1:_(p4) = G_ADDRSPACE_CAST %0:_(pN)
+ //
+ // 3. Global block literals:
+ // %0:_(pN) = G_GLOBAL_VALUE @block_literal_global
+ // %1:_(pN) = G_BITCAST %0 (or spv.bitcast)
+ // %2:_(p4) = G_ADDRSPACE_CAST %1:_(pN)
+ //
+ // 4. Function pointers (direct):
// %0:_(pN) = G_GLOBAL_VALUE @block_invoke_kernel
// %1:_(p4) = G_ADDRSPACE_CAST %0:_(pN)
MachineInstr *MI = MRI->getUniqueVRegDef(ParamReg);
@@ -389,12 +399,30 @@ static MachineInstr *getBlockStructInstr(Register ParamReg,
if (SourceMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
return SourceMI;
- // Otherwise, expect the bitcast sequence (block literal case)
- assert(isSpvIntrinsic(*SourceMI, Intrinsic::spv_bitcast) &&
- SourceMI->getOperand(2).isReg());
- Register ValueReg = SourceMI->getOperand(2).getReg();
- MachineInstr *ValueMI = MRI->getUniqueVRegDef(ValueReg);
- return ValueMI;
+ // Check if it's a direct spv_alloca (stack-allocated block without bitcast)
+ if (isSpvIntrinsic(*SourceMI, Intrinsic::spv_alloca))
+ return SourceMI;
+
+ // Handle G_BITCAST case (stack-allocated blocks)
+ if (SourceMI->getOpcode() == TargetOpcode::G_BITCAST) {
+ assert(SourceMI->getOperand(1).isReg());
+ Register ValueReg = SourceMI->getOperand(1).getReg();
+ MachineInstr *ValueMI = MRI->getUniqueVRegDef(ValueReg);
+ return ValueMI;
+ }
+
+ // Handle spv_bitcast case (global block literal from CrossWorkgroup)
+ if (isSpvIntrinsic(*SourceMI, Intrinsic::spv_bitcast)) {
+ assert(SourceMI->getOperand(2).isReg());
+ Register ValueReg = SourceMI->getOperand(2).getReg();
+ MachineInstr *ValueMI = MRI->getUniqueVRegDef(ValueReg);
+ return ValueMI;
+ }
+
+ // Unhandled case - print debug info
+ SourceMI->print(errs());
+ errs() << "Opcode: " << SourceMI->getOpcode() << "\n";
+ llvm_unreachable("getBlockStructInstr: unexpected instruction pattern");
}
// Return an integer constant corresponding to the given register and
@@ -435,7 +463,8 @@ 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();
+ // return MI->getOperand(1).getGlobal()->getType();
assert(isSpvIntrinsic(*MI, Intrinsic::spv_alloca) &&
"Blocks in OpenCL C must be traceable to allocation site");
return getMachineInstrType(MI);
diff --git a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
index 3fb6810db4e90..5ac36aae452f7 100644
--- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
@@ -1578,6 +1578,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/transcoding/enqueue_kernel.ll b/llvm/test/CodeGen/SPIRV/transcoding/enqueue_kernel.ll
index 616fa6c5fbafb..203f5e6aedd1b 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/enqueue_kernel.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/enqueue_kernel.ll
@@ -2,8 +2,8 @@
; 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: OpCapability DeviceEnqueue
; 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"
@@ -16,17 +16,10 @@
; 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: %[[#EventNull:]] = OpConstantNull %[[#Int8PtrGenTy]]
; CHECK-SPIRV-DAG: %[[#ConstInt21:]] = OpConstant %[[#Int32Ty]] 21{{$}}
; CHECK-SPIRV-DAG: %[[#ConstInt8:]] = OpConstant %[[#Int32Ty]] 8{{$}}
; CHECK-SPIRV-DAG: %[[#ConstInt24:]] = OpConstant %[[#Int32Ty]] 24{{$}}
@@ -46,9 +39,7 @@
;; 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]]
+; CHECK-SPIRV: OpEnqueueKernel %[[#Int32Ty]] %[[#]] %[[#]] %[[#]] %[[#ConstInt0]] %[[#EventNull]] %[[#EventNull]] %[[#BlockKer1]] %[[#]] %[[#ConstInt21]] %[[#ConstInt8]]
;; enqueue_kernel(default_queue, flags, ndrange,
;; ^(void) {
@@ -57,11 +48,7 @@
;; 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]]
+; CHECK-SPIRV: OpEnqueueKernel %[[#Int32Ty]] %[[#]] %[[#]] %[[#]] %[[#ConstInt2]] %[[#]] %[[#]] %[[#BlockKer2]] %[[#]] %[[#ConstInt24]] %[[#ConstInt8]]
;; enqueue_kernel(default_queue, flags, ndrange, 2, &event_wait_list, &clk_event,
;; ^(void) {
@@ -71,12 +58,7 @@
;; 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]]
+; CHECK-SPIRV: OpEnqueueKernel %[[#Int32Ty]] %[[#]] %[[#]] %[[#]] %[[#ConstInt2]] %[[#]] %[[#]] %[[#BlockKer3]] %[[#]] %[[#ConstInt12]] %[[#ConstInt8]]
;; enqueue_kernel(default_queue, flags, ndrange, 2, event_wait_list2, &clk_event,
;; ^(local void *p) {
@@ -86,12 +68,7 @@
;; 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]]
+; CHECK-SPIRV: OpEnqueueKernel %[[#Int32Ty]] %[[#]] %[[#]] %[[#]] %[[#ConstInt0]] %[[#EventNull]] %[[#EventNull]] %[[#BlockKer4]] %[[#]] %[[#ConstInt12]] %[[#ConstInt8]]
;; enqueue_kernel(default_queue, flags, ndrange,
;; ^(local void *p1, local void *p2, local void *p3) {
@@ -101,10 +78,7 @@
;; 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]]
+; CHECK-SPIRV: OpEnqueueKernel %[[#Int32Ty]] %[[#]] %[[#]] %[[#]] %[[#ConstInt0]] %[[#EventNull]] %[[#]] %[[#BlockKer5]] %[[#]] %[[#ConstInt24]] %[[#ConstInt8]]
;; enqueue_kernel(default_queue, flags, ndrange, 0, NULL, &clk_event,
;; ^(void) {
@@ -112,11 +86,11 @@
;; });
;; }
-; 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-SPIRV-DAG: %[[#BlockKer1]] = OpFunction %[[#VoidTy]]
+; CHECK-SPIRV-DAG: %[[#BlockKer2]] = OpFunction %[[#VoidTy]]
+; CHECK-SPIRV-DAG: %[[#BlockKer3]] = OpFunction %[[#VoidTy]]
+; CHECK-SPIRV-DAG: %[[#BlockKer4]] = OpFunction %[[#VoidTy]]
+; CHECK-SPIRV-DAG: %[[#BlockKer5]] = OpFunction %[[#VoidTy]]
%opencl.queue_t = type opaque
%struct.ndrange_t = type { i32 }
>From 8a83da9752f31a97346b706ffec93f1c4bdd79bb Mon Sep 17 00:00:00 2001
From: idubinov <igor.dubinov at amd.com>
Date: Fri, 20 Mar 2026 06:37:22 -0500
Subject: [PATCH 04/12] Revert "fix enqueue kernel test"
This reverts commit 1faa1e0a54c8a90f7dd933ebf9ecce9fb6e797a2.
---
llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp | 49 ++++--------------
llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp | 1 -
.../SPIRV/transcoding/enqueue_kernel.ll | 50 ++++++++++++++-----
3 files changed, 48 insertions(+), 52 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
index 101222651e330..bb88de34aa7f0 100644
--- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
@@ -370,23 +370,13 @@ lookupBuiltin(StringRef DemangledCall,
static MachineInstr *getBlockStructInstr(Register ParamReg,
MachineRegisterInfo *MRI) {
- // We expect one of the following sequences of instructions:
- //
- // 1. Stack-allocated blocks (with bitcast):
+ // 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)
//
- // 2. Stack-allocated blocks (direct):
- // %0:_(pN) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.alloca)
- // %1:_(p4) = G_ADDRSPACE_CAST %0:_(pN)
- //
- // 3. Global block literals:
- // %0:_(pN) = G_GLOBAL_VALUE @block_literal_global
- // %1:_(pN) = G_BITCAST %0 (or spv.bitcast)
- // %2:_(p4) = G_ADDRSPACE_CAST %1:_(pN)
- //
- // 4. Function pointers (direct):
+ // For function pointers (empty kernel blocks), the sequence may be simpler:
// %0:_(pN) = G_GLOBAL_VALUE @block_invoke_kernel
// %1:_(p4) = G_ADDRSPACE_CAST %0:_(pN)
MachineInstr *MI = MRI->getUniqueVRegDef(ParamReg);
@@ -399,30 +389,12 @@ static MachineInstr *getBlockStructInstr(Register ParamReg,
if (SourceMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
return SourceMI;
- // Check if it's a direct spv_alloca (stack-allocated block without bitcast)
- if (isSpvIntrinsic(*SourceMI, Intrinsic::spv_alloca))
- return SourceMI;
-
- // Handle G_BITCAST case (stack-allocated blocks)
- if (SourceMI->getOpcode() == TargetOpcode::G_BITCAST) {
- assert(SourceMI->getOperand(1).isReg());
- Register ValueReg = SourceMI->getOperand(1).getReg();
- MachineInstr *ValueMI = MRI->getUniqueVRegDef(ValueReg);
- return ValueMI;
- }
-
- // Handle spv_bitcast case (global block literal from CrossWorkgroup)
- if (isSpvIntrinsic(*SourceMI, Intrinsic::spv_bitcast)) {
- assert(SourceMI->getOperand(2).isReg());
- Register ValueReg = SourceMI->getOperand(2).getReg();
- MachineInstr *ValueMI = MRI->getUniqueVRegDef(ValueReg);
- return ValueMI;
- }
-
- // Unhandled case - print debug info
- SourceMI->print(errs());
- errs() << "Opcode: " << SourceMI->getOpcode() << "\n";
- llvm_unreachable("getBlockStructInstr: unexpected instruction pattern");
+ // Otherwise, expect the bitcast sequence (block literal case)
+ assert(isSpvIntrinsic(*SourceMI, Intrinsic::spv_bitcast) &&
+ SourceMI->getOperand(2).isReg());
+ Register ValueReg = SourceMI->getOperand(2).getReg();
+ MachineInstr *ValueMI = MRI->getUniqueVRegDef(ValueReg);
+ return ValueMI;
}
// Return an integer constant corresponding to the given register and
@@ -463,8 +435,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()->getValueType();
- // return MI->getOperand(1).getGlobal()->getType();
+ return MI->getOperand(1).getGlobal()->getType();
assert(isSpvIntrinsic(*MI, Intrinsic::spv_alloca) &&
"Blocks in OpenCL C must be traceable to allocation site");
return getMachineInstrType(MI);
diff --git a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
index 5ac36aae452f7..3fb6810db4e90 100644
--- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
@@ -1578,7 +1578,6 @@ 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/transcoding/enqueue_kernel.ll b/llvm/test/CodeGen/SPIRV/transcoding/enqueue_kernel.ll
index 203f5e6aedd1b..616fa6c5fbafb 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/enqueue_kernel.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/enqueue_kernel.ll
@@ -2,8 +2,8 @@
; 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: OpCapability DeviceEnqueue
; 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"
@@ -16,10 +16,17 @@
; 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 %[[#Int8PtrGenTy]]
+; 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{{$}}
@@ -39,7 +46,9 @@
;; Emits block literal on stack and block kernel.
-; CHECK-SPIRV: OpEnqueueKernel %[[#Int32Ty]] %[[#]] %[[#]] %[[#]] %[[#ConstInt0]] %[[#EventNull]] %[[#EventNull]] %[[#BlockKer1]] %[[#]] %[[#ConstInt21]] %[[#ConstInt8]]
+; 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) {
@@ -48,7 +57,11 @@
;; Emits block literal on stack and block kernel.
-; CHECK-SPIRV: OpEnqueueKernel %[[#Int32Ty]] %[[#]] %[[#]] %[[#]] %[[#ConstInt2]] %[[#]] %[[#]] %[[#BlockKer2]] %[[#]] %[[#ConstInt24]] %[[#ConstInt8]]
+; 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) {
@@ -58,7 +71,12 @@
;; char c;
;; Emits global block literal and block kernel.
-; CHECK-SPIRV: OpEnqueueKernel %[[#Int32Ty]] %[[#]] %[[#]] %[[#]] %[[#ConstInt2]] %[[#]] %[[#]] %[[#BlockKer3]] %[[#]] %[[#ConstInt12]] %[[#ConstInt8]]
+; 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) {
@@ -68,7 +86,12 @@
;; Emits global block literal and block kernel.
-; CHECK-SPIRV: OpEnqueueKernel %[[#Int32Ty]] %[[#]] %[[#]] %[[#]] %[[#ConstInt0]] %[[#EventNull]] %[[#EventNull]] %[[#BlockKer4]] %[[#]] %[[#ConstInt12]] %[[#ConstInt8]]
+; 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) {
@@ -78,7 +101,10 @@
;; Emits block literal on stack and block kernel.
-; CHECK-SPIRV: OpEnqueueKernel %[[#Int32Ty]] %[[#]] %[[#]] %[[#]] %[[#ConstInt0]] %[[#EventNull]] %[[#]] %[[#BlockKer5]] %[[#]] %[[#ConstInt24]] %[[#ConstInt8]]
+; 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) {
@@ -86,11 +112,11 @@
;; });
;; }
-; CHECK-SPIRV-DAG: %[[#BlockKer1]] = OpFunction %[[#VoidTy]]
-; CHECK-SPIRV-DAG: %[[#BlockKer2]] = OpFunction %[[#VoidTy]]
-; CHECK-SPIRV-DAG: %[[#BlockKer3]] = OpFunction %[[#VoidTy]]
-; CHECK-SPIRV-DAG: %[[#BlockKer4]] = OpFunction %[[#VoidTy]]
-; CHECK-SPIRV-DAG: %[[#BlockKer5]] = OpFunction %[[#VoidTy]]
+; 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]]
%opencl.queue_t = type opaque
%struct.ndrange_t = type { i32 }
>From ed2152762132724a49916d76833351e4bc877f9d Mon Sep 17 00:00:00 2001
From: idubinov <igor.dubinov at amd.com>
Date: Fri, 20 Mar 2026 10:19:01 -0500
Subject: [PATCH 05/12] add test no events
---
llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll | 39 +++++++++++++++++++
1 file changed, 39 insertions(+)
diff --git a/llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll b/llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll
index 442f0dc3cf486..f8f41e3488c84 100644
--- a/llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll
+++ b/llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll
@@ -64,3 +64,42 @@ declare i32 @__enqueue_kernel_basic_events(ptr, i32, ptr, i32, ptr addrspace(4),
; CHECK-SPIRV: %[[#Invoke]] = OpFunction %[[#Void]] None %[[#]]
; CHECK-SPIRV-NEXT: %[[#]] = OpFunctionParameter %[[#Int8PtrGen]]
+
+;; Test case for enqueueing empty kernel with no events
+;; __kernel void test_enqueue_empty_no_events() {
+;; enqueue_kernel(get_default_queue(),
+;; CLK_ENQUEUE_FLAGS_WAIT_KERNEL,
+;; ndrange_1D(1),
+;; ^(){});
+;; }
+
+ at __block_literal_global_no_events = internal addrspace(1) constant { i32, i32 } { i32 8, i32 4 }, align 4
+
+define spir_kernel void @test_enqueue_empty_no_events() {
+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(ptr %call, i32 1, ptr %tmp, ptr addrspace(4) addrspacecast (ptr @__test_enqueue_empty_no_events_block_invoke_kernel to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @__block_literal_global_no_events to ptr addrspace(4)))
+ ret void
+; CHECK-SPIRV: %[[#]] = OpEnqueueKernel %[[#]] %[[#]] %[[#]] %[[#]] %[[#Invoke2:]] %[[#]] %[[#]] %[[#]]
+}
+
+define internal spir_func void @__test_enqueue_empty_no_events_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_no_events_block_invoke_kernel(ptr addrspace(4)) {
+entry:
+ call void @__test_enqueue_empty_no_events_block_invoke(ptr addrspace(4) %0)
+ ret void
+}
+
+declare i32 @__enqueue_kernel_basic(ptr, i32, ptr, ptr addrspace(4), ptr addrspace(4))
+
+; CHECK-SPIRV-DAG: %[[#Invoke2:]] = OpFunction %[[#Void]] None %[[#]]
+; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#Int8PtrGen]]
>From df368f99c5f5a15b264aad45034fc874bec56318 Mon Sep 17 00:00:00 2001
From: idubinov <igor.dubinov at amd.com>
Date: Fri, 20 Mar 2026 10:19:32 -0500
Subject: [PATCH 06/12] remove todo
---
llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll | 3 ---
1 file changed, 3 deletions(-)
diff --git a/llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll b/llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll
index f8f41e3488c84..d54283d221883 100644
--- a/llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll
+++ b/llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll
@@ -1,8 +1,5 @@
; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-; TODO(#60133): Requires updates following opaque pointer migration.
-
-
;; 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:
>From 39f8ae6081917dc405481a5bfe6c186062db11db Mon Sep 17 00:00:00 2001
From: idubinov <idubinov at amd.com>
Date: Mon, 23 Mar 2026 11:43:31 +0100
Subject: [PATCH 07/12] Update llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
Co-authored-by: Arseniy Obolenskiy <gooddoog at student.su>
---
llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp | 1 +
1 file changed, 1 insertion(+)
diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
index bb88de34aa7f0..0440b98bb98cb 100644
--- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
@@ -384,6 +384,7 @@ static MachineInstr *getBlockStructInstr(Register ParamReg,
MI->getOperand(1).isReg());
Register CastSourceReg = MI->getOperand(1).getReg();
MachineInstr *SourceMI = MRI->getUniqueVRegDef(CastSourceReg);
+ assert(SourceMI);
// Check if it's a direct G_GLOBAL_VALUE (function pointer case)
if (SourceMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
>From 8f7570cce7851f79010c05e291a954d593bb1e2d Mon Sep 17 00:00:00 2001
From: idubinov <idubinov at amd.com>
Date: Mon, 23 Mar 2026 12:10:12 +0100
Subject: [PATCH 08/12] Apply suggestions from code review
Co-authored-by: Marcos Maronas <mmaronas at amd.com>
---
llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
index 0440b98bb98cb..e1792844179fe 100644
--- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
@@ -386,11 +386,11 @@ static MachineInstr *getBlockStructInstr(Register ParamReg,
MachineInstr *SourceMI = MRI->getUniqueVRegDef(CastSourceReg);
assert(SourceMI);
- // Check if it's a direct G_GLOBAL_VALUE (function pointer case)
+ // Check if it's a direct G_GLOBAL_VALUE (function pointer case).
if (SourceMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
return SourceMI;
- // Otherwise, expect the bitcast sequence (block literal case)
+ // Otherwise, expect the bitcast sequence (block literal case).
assert(isSpvIntrinsic(*SourceMI, Intrinsic::spv_bitcast) &&
SourceMI->getOperand(2).isReg());
Register ValueReg = SourceMI->getOperand(2).getReg();
>From 024c92fa617ecefb8fc3a77db1e1d3728a2b27cf Mon Sep 17 00:00:00 2001
From: idubinov <idubinov at amd.com>
Date: Mon, 23 Mar 2026 12:48:58 +0100
Subject: [PATCH 09/12] Update llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
Co-authored-by: Dmitry Sidorov <dsidorov at amd.com>
---
llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
index e1792844179fe..96f4ea3ed8f20 100644
--- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
@@ -2757,7 +2757,7 @@ static bool buildEnqueueKernel(const SPIRV::IncomingCall *Call,
Register BlockLiteralReg = Call->Arguments[BlockFIdx + 1];
Type *PType = const_cast<Type *>(getBlockStructType(BlockLiteralReg, MRI));
- // OpEnqueueKernel requires the Param to be a pointer to i8 (per SPIR-V spec).
+ // OpEnqueueKernel requires the Param to be a pointer to i8.
// BlockLiteralReg is a Generic pointer to the block struct.
// Bitcast it to a Generic pointer to i8.
const SPIRVTypeInst Int8Ty = GR->getOrCreateSPIRVIntegerType(8, MIRBuilder);
>From de22854d19749ae15cd67a8d23f4194051453274 Mon Sep 17 00:00:00 2001
From: idubinov <igor.dubinov at amd.com>
Date: Mon, 23 Mar 2026 07:56:15 -0500
Subject: [PATCH 10/12] address review
---
llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp | 8 ++++----
llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll | 3 ++-
2 files changed, 6 insertions(+), 5 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
index 96f4ea3ed8f20..31149c431bdd2 100644
--- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
@@ -384,7 +384,7 @@ static MachineInstr *getBlockStructInstr(Register ParamReg,
MI->getOperand(1).isReg());
Register CastSourceReg = MI->getOperand(1).getReg();
MachineInstr *SourceMI = MRI->getUniqueVRegDef(CastSourceReg);
- assert(SourceMI);
+ assert(SourceMI && "Definition for source reg not found.");
// Check if it's a direct G_GLOBAL_VALUE (function pointer case).
if (SourceMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
@@ -2755,7 +2755,7 @@ static bool buildEnqueueKernel(const SPIRV::IncomingCall *Call,
assert(BlockMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE);
Register BlockLiteralReg = Call->Arguments[BlockFIdx + 1];
- Type *PType = const_cast<Type *>(getBlockStructType(BlockLiteralReg, MRI));
+ const Type *PType = getBlockStructType(BlockLiteralReg, MRI);
// OpEnqueueKernel requires the Param to be a pointer to i8.
// BlockLiteralReg is a Generic pointer to the block struct.
@@ -2797,9 +2797,9 @@ static bool buildEnqueueKernel(const SPIRV::IncomingCall *Call,
// 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));
+ MIB.addUse(buildConstantIntReg32(DL.getTypeStoreSize(const_cast<Type *>(PType)), MIRBuilder, GR));
// Param Aligment: Aligment of block literal structure.
- MIB.addUse(buildConstantIntReg32(DL.getPrefTypeAlign(PType).value(),
+ MIB.addUse(buildConstantIntReg32(DL.getPrefTypeAlign(const_cast<Type *>(PType)).value(),
MIRBuilder, GR));
for (unsigned i = 0; i < LocalSizes.size(); i++)
diff --git a/llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll b/llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll
index d54283d221883..4c038e372e16a 100644
--- a/llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll
+++ b/llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll
@@ -1,4 +1,5 @@
-; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: llc -O0 -verify-machineinstrs -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: %if spirv-tools %{ llc -O0 -verify-machineinstrs -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
;; This test checks that Invoke parameter of OpEnueueKernel instruction meet the
;; following specification requirements in case of enqueueing empty block:
>From d6db0059ba5cbc1dfa1d6c71078f94a4a8bf7ff6 Mon Sep 17 00:00:00 2001
From: idubinov <igor.dubinov at amd.com>
Date: Tue, 7 Apr 2026 11:00:31 -0500
Subject: [PATCH 11/12] address review
---
llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp | 225 +++---
llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp | 1 +
llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll | 103 ---
.../SPIRV/transcoding/enqueue_kernel.ll | 665 +++++++++---------
4 files changed, 472 insertions(+), 522 deletions(-)
delete mode 100644 llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll
diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
index 31149c431bdd2..95f4466ad94e3 100644
--- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
@@ -370,13 +370,23 @@ lookupBuiltin(StringRef DemangledCall,
static MachineInstr *getBlockStructInstr(Register ParamReg,
MachineRegisterInfo *MRI) {
- // We expect the following sequence of instructions:
+ // We expect one of the following sequences of instructions:
+ //
+ // 1. Stack-allocated blocks (with bitcast):
// %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)
//
- // For function pointers (empty kernel blocks), the sequence may be simpler:
+ // 2. Stack-allocated blocks (direct):
+ // %0:_(pN) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.spv.alloca)
+ // %1:_(p4) = G_ADDRSPACE_CAST %0:_(pN)
+ //
+ // 3. Global block literals:
+ // %0:_(pN) = G_GLOBAL_VALUE @block_literal_global
+ // %1:_(pN) = G_BITCAST %0 (or spv.bitcast)
+ // %2:_(p4) = G_ADDRSPACE_CAST %1:_(pN)
+ //
+ // 4. Function pointers (direct):
// %0:_(pN) = G_GLOBAL_VALUE @block_invoke_kernel
// %1:_(p4) = G_ADDRSPACE_CAST %0:_(pN)
MachineInstr *MI = MRI->getUniqueVRegDef(ParamReg);
@@ -386,16 +396,34 @@ static MachineInstr *getBlockStructInstr(Register ParamReg,
MachineInstr *SourceMI = MRI->getUniqueVRegDef(CastSourceReg);
assert(SourceMI && "Definition for source reg not found.");
- // Check if it's a direct G_GLOBAL_VALUE (function pointer case).
+ // Check if it's a direct G_GLOBAL_VALUE (function pointer case)
if (SourceMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
return SourceMI;
- // Otherwise, expect the bitcast sequence (block literal case).
- assert(isSpvIntrinsic(*SourceMI, Intrinsic::spv_bitcast) &&
- SourceMI->getOperand(2).isReg());
- Register ValueReg = SourceMI->getOperand(2).getReg();
- MachineInstr *ValueMI = MRI->getUniqueVRegDef(ValueReg);
- return ValueMI;
+ // Check if it's a direct spv_alloca (stack-allocated block without bitcast)
+ if (isSpvIntrinsic(*SourceMI, Intrinsic::spv_alloca))
+ return SourceMI;
+
+ // Handle G_BITCAST case (stack-allocated blocks)
+ if (SourceMI->getOpcode() == TargetOpcode::G_BITCAST) {
+ assert(SourceMI->getOperand(1).isReg());
+ Register ValueReg = SourceMI->getOperand(1).getReg();
+ MachineInstr *ValueMI = MRI->getUniqueVRegDef(ValueReg);
+ return ValueMI;
+ }
+
+ // Handle spv_bitcast case (global block literal from CrossWorkgroup)
+ if (isSpvIntrinsic(*SourceMI, Intrinsic::spv_bitcast)) {
+ assert(SourceMI->getNumOperands()>=2 && SourceMI->getOperand(2).isReg());
+ Register ValueReg = SourceMI->getOperand(2).getReg();
+ MachineInstr *ValueMI = MRI->getUniqueVRegDef(ValueReg);
+ return ValueMI;
+ }
+
+ // Unhandled case - print debug info
+ SourceMI->print(errs());
+ errs() << "Opcode: " << SourceMI->getOpcode() << "\n";
+ llvm_unreachable("getBlockStructInstr: unexpected instruction pattern");
}
// Return an integer constant corresponding to the given register and
@@ -2707,103 +2735,130 @@ getOrCreateSPIRVDeviceEventPointer(MachineIRBuilder &MIRBuilder,
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.
- 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) {
- 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.
- .addUse(buildConstantIntReg32(0, MIRBuilder, GR)) // Indices.
- .addUse(buildConstantIntReg32(I, MIRBuilder, GR));
- LocalSizes.push_back(Reg);
- }
- }
+ // 1. prepare call indexes in order we expect them.
+ // Based on clang sources, clang/lib/CodeGen/CGBuiltin.cpp, BIenqueue_kernel,
+ // We expect 4 different 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
+ /// TODO: handle __spirv_EnqueueKernel
- // Prepare block invoke function and block literal before building
- // OpEnqueueKernel.
- const unsigned BlockFIdx = HasEvents ? 6 : 3;
- MachineInstr *BlockMI = getBlockStructInstr(Call->Arguments[BlockFIdx], MRI);
- assert(BlockMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE);
+ 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);
+ Register NullPtr = GR->getOrCreateConstNullPtr(
+ MIRBuilder, getOrCreateSPIRVDeviceEventPointer(MIRBuilder, GR));
+ WaitEventsReg = NullPtr;
+ RetEventReg = NullPtr;
+ }
- Register BlockLiteralReg = Call->Arguments[BlockFIdx + 1];
- const Type *PType = getBlockStructType(BlockLiteralReg, MRI);
+ /// 2.2 Invoke (Kernel)
+ assert(getBlockStructInstr(Call->Arguments[InvokeIdx], MRI)->getOpcode() == TargetOpcode::G_GLOBAL_VALUE);
- // OpEnqueueKernel requires the Param to be a pointer to i8.
- // BlockLiteralReg is a Generic pointer to the block struct.
- // Bitcast it to a Generic pointer to i8.
+ /// 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);
+ const Type *PType = getBlockStructType(BlockLiteralReg, MRI);
- Register BlockLiteralGenAsI8 =
+ /// TODO: check if already Int8Ptr
+ Register ParamReg =
createVirtualRegister(Int8PtrGen, GR, MIRBuilder);
MIRBuilder.buildInstr(SPIRV::OpBitcast)
- .addDef(BlockLiteralGenAsI8)
+ .addDef(ParamReg)
.addUse(GR->getSPIRVTypeID(Int8PtrGen))
.addUse(BlockLiteralReg);
+ /// TODO: these numbers should be obtained from block literal structure.
+ Register ParamSizeReg = buildConstantIntReg32(DL.getTypeStoreSize(const_cast<Type *>(PType)), MIRBuilder, GR);
+ Register ParamAlignReg = buildConstantIntReg32(DL.getPrefTypeAlign(const_cast<Type *>(PType)).value(), MIRBuilder, GR);
- // SPIRV OpEnqueueKernel instruction has 10+ arguments.
- auto MIB = MIRBuilder.buildInstr(SPIRV::OpEnqueueKernel)
- .addDef(Call->ReturnRegister)
- .addUse(GR->getSPIRVTypeID(Int32Ty));
- // Copy all arguments before block invoke function pointer.
- for (unsigned i = 0; i < BlockFIdx; i++)
- MIB.addUse(Call->Arguments[i]);
+ /// 2.4 Local Size Array
+ SmallVector<Register, 16> LocalSizes;
+ 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();
- // 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.
- }
+ Register LocalSizeArrayReg = Call->Arguments[LocalSizeElemPtrIdx];
- // Invoke: Pointer to invoke function.
- MIB.addGlobalAddress(BlockMI->getOperand(1).getGlobal());
+ for (unsigned i = 0; i < NumElem; ++i) {
+ Register Reg = MRI->createVirtualRegister(&SPIRV::pIDRegClass);
+ auto GEPInst = MIRBuilder.buildIntrinsic(
+ Intrinsic::spv_gep, ArrayRef<Register>{Reg}, true, false);
+ GEPInst
+ .addImm(0) // In bound.
+ .addUse(LocalSizeArrayReg) // Base pointer.
+ .addUse(buildConstantIntReg32(0, MIRBuilder, GR)) // Indices.
+ .addUse(buildConstantIntReg32(i, MIRBuilder, GR));
+ LocalSizes.push_back(Reg);
+ }
+ }
- // Param: Pointer to block literal (as Generic i8*).
- MIB.addUse(BlockLiteralGenAsI8);
- // TODO: these numbers should be obtained from block literal structure.
- // Param Size: Size of block literal structure.
- MIB.addUse(buildConstantIntReg32(DL.getTypeStoreSize(const_cast<Type *>(PType)), MIRBuilder, GR));
- // Param Aligment: Aligment of block literal structure.
- MIB.addUse(buildConstantIntReg32(DL.getPrefTypeAlign(const_cast<Type *>(PType)).value(),
- MIRBuilder, GR));
+ /// 3. create a SPIRV operator with arguments.
+ auto MIB = MIRBuilder.buildInstr(SPIRV::OpEnqueueKernel)
+ .addDef(Call->ReturnRegister)
+ .addUse(GR->getSPIRVTypeID(Int32Ty))
+ .addUse(Call->Arguments[QueueIdx])
+ .addUse(Call->Arguments[FlagsIdx])
+ .addUse(Call->Arguments[NDRangeIdx])
+ .addUse(NumEventsReg)
+ .addUse(WaitEventsReg)
+ .addUse(RetEventReg)
+ .addUse(Call->Arguments[InvokeIdx])
+ .addUse(ParamReg)
+ .addUse(ParamSizeReg)
+ .addUse(ParamAlignReg);
+ for (auto & LocalSize: LocalSizes)
+ MIB.addUse(LocalSize);
- for (unsigned i = 0; i < LocalSizes.size(); i++)
- MIB.addUse(LocalSizes[i]);
return true;
}
diff --git a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
index 3fb6810db4e90..5ac36aae452f7 100644
--- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
@@ -1578,6 +1578,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 4c038e372e16a..0000000000000
--- a/llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll
+++ /dev/null
@@ -1,103 +0,0 @@
-; RUN: llc -O0 -verify-machineinstrs -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-; RUN: %if spirv-tools %{ llc -O0 -verify-machineinstrs -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
-
-;; 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: %[[#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: %[[#BlockGeneric:]] = OpPtrCastToGeneric %[[#]] %[[#Block]]
-; CHECK-SPIRV: %[[#Int8PtrGenBlock:]] = OpBitcast %[[#Int8PtrGen]] %[[#BlockGeneric]]
-; 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]]
-
-;; Test case for enqueueing empty kernel with no events
-;; __kernel void test_enqueue_empty_no_events() {
-;; enqueue_kernel(get_default_queue(),
-;; CLK_ENQUEUE_FLAGS_WAIT_KERNEL,
-;; ndrange_1D(1),
-;; ^(){});
-;; }
-
- at __block_literal_global_no_events = internal addrspace(1) constant { i32, i32 } { i32 8, i32 4 }, align 4
-
-define spir_kernel void @test_enqueue_empty_no_events() {
-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(ptr %call, i32 1, ptr %tmp, ptr addrspace(4) addrspacecast (ptr @__test_enqueue_empty_no_events_block_invoke_kernel to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @__block_literal_global_no_events to ptr addrspace(4)))
- ret void
-; CHECK-SPIRV: %[[#]] = OpEnqueueKernel %[[#]] %[[#]] %[[#]] %[[#]] %[[#Invoke2:]] %[[#]] %[[#]] %[[#]]
-}
-
-define internal spir_func void @__test_enqueue_empty_no_events_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_no_events_block_invoke_kernel(ptr addrspace(4)) {
-entry:
- call void @__test_enqueue_empty_no_events_block_invoke(ptr addrspace(4) %0)
- ret void
-}
-
-declare i32 @__enqueue_kernel_basic(ptr, i32, ptr, ptr addrspace(4), ptr addrspace(4))
-
-; CHECK-SPIRV-DAG: %[[#Invoke2:]] = OpFunction %[[#Void]] None %[[#]]
-; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#Int8PtrGen]]
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/enqueue_kernel.ll b/llvm/test/CodeGen/SPIRV/transcoding/enqueue_kernel.ll
index 616fa6c5fbafb..a0274b9ad0205 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/enqueue_kernel.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/enqueue_kernel.ll
@@ -1,389 +1,386 @@
-; 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: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val
+; 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: %[[#Num29i32:]] = OpConstant %[[#typeInt32]] 29
+; CHECK-DAG: %[[#Num36i32:]] = OpConstant %[[#typeInt32]] 36
+
+; 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-LABEL: ; -- Begin function device_side_enqueue
+
+; CHECK: %[[#NDRange3sret:]] = OpBuildNDRange %[[#TypeNDRangeStruct]] %[[#]] %[[#nullArray3x64]] %[[#nullArray3x64]]
+; 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};
+;;
+;; // enqueue empty kernel
+; CHECK: %[[#]] = OpEnqueueKernel %[[#typeInt32]] %[[#default_queue:]] %[[#Num1i32]] %[[#NDRange3]] %[[#Num0i32]] %[[#nullPtrInt8]] %[[#nullPtrInt8]] %[[#]] %[[#]] %[[#Num8i32]] %[[#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]] %[[#nullPtrInt8]] %[[#nullPtrInt8]] %[[#]] %[[#]] %[[#Num29i32]] %[[#Num8i32]]
+;; enqueue_kernel(default_queue, flags, ndrange,
+;; ^(void) {
+;; a[i] = c0;
+;; });
+;;
+;; // event, no var args
+; CHECK: %[[#]] = OpEnqueueKernel %[[#typeInt32]] %[[#default_queue]] %[[#Num0i32]] %[[#]] %[[#Num2i32]] %[[#event_wait_list:]] %[[#clk_event:]] %[[#]] %[[#]] %[[#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:]] %[[#clk_event]] %[[#]] %[[#]] %[[#Num8i32]] %[[#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]] %[[#nullPtrInt8]] %[[#nullPtrInt8]] %[[#]] %[[#]] %[[#Num8i32]] %[[#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]] %[[#clk_event]] %[[#]] %[[#]] %[[#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]]
-%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 __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 #0
+ 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 #0
+ 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 #0
- 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
+; Function Attrs: convergent norecurse nounwind
+define spir_kernel void @device_side_enqueue(ptr addrspace(1) noundef align 4 %a, ptr addrspace(1) noundef align 4 %b, i32 noundef %i, i8 noundef signext %c0) local_unnamed_addr #1 !kernel_arg_addr_space !6 !kernel_arg_access_qual !7 !kernel_arg_type !8 !kernel_arg_base_type !8 !kernel_arg_type_qual !9 {
+entry:
+ %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.lifetime.start.p0(ptr nonnull %tmp.i)
+ call void @llvm.lifetime.start.p0(ptr nonnull %tmp1.i)
+ call void @llvm.lifetime.start.p0(ptr nonnull %block.i)
+ call void @llvm.lifetime.start.p0(ptr nonnull %tmp4.i)
+ call void @llvm.lifetime.start.p0(ptr nonnull %block5.i)
+ call void @llvm.lifetime.start.p0(ptr nonnull %tmp12.i)
+ call void @llvm.lifetime.start.p0(ptr nonnull %tmp14.i)
+ call void @llvm.lifetime.start.p0(ptr nonnull %tmp16.i)
+ call void @llvm.lifetime.start.p0(ptr nonnull %block17.i)
+ call void @llvm.lifetime.start.p0(ptr nonnull %clk_event.i) #8
+ call void @llvm.lifetime.start.p0(ptr nonnull %event_wait_list.i) #8
+ call void @llvm.lifetime.start.p0(ptr nonnull %event_wait_list2.i) #8
+ call void @llvm.lifetime.start.p0(ptr nonnull %gs.i) #8
+ call void @llvm.memcpy.p0.p2.i64(ptr noundef nonnull align 8 dereferenceable(24) %gs.i, ptr addrspace(2) noundef align 8 dereferenceable(24) @__const.device_side_enqueue.gs, i64 24, i1 false)
+ call spir_func void @_Z10ndrange_3DPKm(ptr dead_on_unwind nonnull writable sret(%struct.ndrange_t) align 8 %tmp.i, ptr noundef nonnull %gs.i) #9
+ %0 = call spir_func i32 @__enqueue_kernel_basic_events(target("spirv.Queue") undef, i32 1, ptr nonnull %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))) #8
+ store i32 29, ptr %block.i, align 8
+ %block.align.i = getelementptr inbounds nuw i8, ptr %block.i, i64 4
+ store i32 8, ptr %block.align.i, align 4
+ %block.invoke.i = getelementptr inbounds nuw 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 nuw i8, ptr %block.i, i64 16
+ store ptr addrspace(1) %a, ptr %block.captured.i, align 8, !tbaa !10
+ %block.captured2.i = getelementptr inbounds nuw i8, ptr %block.i, i64 24
+ store i32 %i, ptr %block.captured2.i, align 8, !tbaa !2
+ %block.captured3.i = getelementptr inbounds nuw i8, ptr %block.i, i64 28
+ store i8 %c0, ptr %block.captured3.i, align 4, !tbaa !13
+ %1 = addrspacecast ptr %block.i to ptr addrspace(4)
+ %2 = call spir_func i32 @__enqueue_kernel_basic(target("spirv.Queue") undef, i32 0, ptr nonnull %tmp1.i, ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_2_kernel to ptr addrspace(4)), ptr addrspace(4) %1) #8
+ %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 nuw i8, ptr %block5.i, i64 4
+ store i32 8, ptr %block.align7.i, align 4
+ %block.invoke8.i = getelementptr inbounds nuw 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 nuw i8, ptr %block5.i, i64 16
+ store ptr addrspace(1) %a, ptr %block.captured9.i, align 8, !tbaa !10
+ %block.captured10.i = getelementptr inbounds nuw i8, ptr %block5.i, i64 32
+ store i32 %i, ptr %block.captured10.i, align 8, !tbaa !2
+ %block.captured11.i = getelementptr inbounds nuw i8, ptr %block5.i, i64 24
+ store ptr addrspace(1) %b, ptr %block.captured11.i, align 8, !tbaa !10
+ %5 = addrspacecast ptr %block5.i to ptr addrspace(4)
+ %6 = call spir_func i32 @__enqueue_kernel_basic_events(target("spirv.Queue") undef, i32 0, ptr nonnull %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) #8
+ %7 = addrspacecast ptr %event_wait_list2.i to ptr addrspace(4)
+ call void @llvm.lifetime.start.p0(ptr nonnull %block_sizes.i) #8
+ store i64 0, ptr %block_sizes.i, align 8
+ %8 = call spir_func i32 @__enqueue_kernel_events_varargs(target("spirv.Queue") undef, i32 0, ptr nonnull %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 nonnull %block_sizes.i) #8
+ call void @llvm.lifetime.end.p0(ptr nonnull %block_sizes.i) #8
+ call void @llvm.lifetime.start.p0(ptr nonnull %block_sizes15.i) #8
+ store i64 101, ptr %block_sizes15.i, align 8
+ %9 = getelementptr inbounds nuw i8, ptr %block_sizes15.i, i64 8
+ store i64 102, ptr %9, align 8
+ %10 = getelementptr inbounds nuw 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") undef, i32 0, ptr nonnull %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 nonnull %block_sizes15.i) #8
+ call void @llvm.lifetime.end.p0(ptr nonnull %block_sizes15.i) #8
+ store i32 36, ptr %block17.i, align 8
+ %block.align19.i = getelementptr inbounds nuw i8, ptr %block17.i, i64 4
+ store i32 8, ptr %block.align19.i, align 4
+ %block.invoke20.i = getelementptr inbounds nuw 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 nuw i8, ptr %block17.i, i64 16
+ store ptr addrspace(1) %a, ptr %block.captured21.i, align 8, !tbaa !10
+ %block.captured22.i = getelementptr inbounds nuw i8, ptr %block17.i, i64 32
+ store i32 %i, ptr %block.captured22.i, align 8, !tbaa !2
+ %block.captured23.i = getelementptr inbounds nuw i8, ptr %block17.i, i64 24
+ store ptr addrspace(1) %b, ptr %block.captured23.i, align 8, !tbaa !10
+ %12 = addrspacecast ptr %block17.i to ptr addrspace(4)
+ %13 = call spir_func i32 @__enqueue_kernel_basic_events(target("spirv.Queue") undef, i32 0, ptr nonnull %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) #8
+ call void @llvm.lifetime.end.p0(ptr nonnull %gs.i) #8
+ call void @llvm.lifetime.end.p0(ptr nonnull %event_wait_list2.i) #8
+ call void @llvm.lifetime.end.p0(ptr nonnull %event_wait_list.i) #8
+ call void @llvm.lifetime.end.p0(ptr nonnull %clk_event.i) #8
+ call void @llvm.lifetime.end.p0(ptr nonnull %tmp.i)
+ call void @llvm.lifetime.end.p0(ptr nonnull %tmp1.i)
+ call void @llvm.lifetime.end.p0(ptr nonnull %block.i)
+ call void @llvm.lifetime.end.p0(ptr nonnull %tmp4.i)
+ call void @llvm.lifetime.end.p0(ptr nonnull %block5.i)
+ call void @llvm.lifetime.end.p0(ptr nonnull %tmp12.i)
+ call void @llvm.lifetime.end.p0(ptr nonnull %tmp14.i)
+ call void @llvm.lifetime.end.p0(ptr nonnull %tmp16.i)
+ call void @llvm.lifetime.end.p0(ptr nonnull %block17.i)
+ ret void
+}
+
+; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite)
+declare void @llvm.lifetime.start.p0(ptr captures(none)) #2
-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) {
+; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: readwrite)
+declare void @llvm.memcpy.p0.p2.i64(ptr noalias writeonly captures(none), ptr addrspace(2) noalias readonly captures(none), i64, i1 immarg) #3
+
+; Function Attrs: convergent nounwind
+declare spir_func void @_Z10ndrange_3DPKm(ptr dead_on_unwind writable sret(%struct.ndrange_t) align 8, ptr noundef) local_unnamed_addr #4
+
+; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(none)
+define internal spir_func void @__device_side_enqueue_block_invoke(ptr addrspace(4) readnone captures(none) %.block_descriptor) #5 {
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)
ret void
}
-declare void @llvm.memcpy.p0.p0.i32(ptr noalias nocapture writeonly, ptr noalias nocapture readonly, i32, i1 immarg)
+; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(none)
+define internal spir_kernel void @__device_side_enqueue_block_invoke_kernel(ptr addrspace(4) readnone captures(none) %0) #5 {
+entry:
+ ret void
+}
+
+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)) local_unnamed_addr
-define internal spir_func void @__device_side_enqueue_block_invoke(ptr addrspace(4) noundef %.block_descriptor) {
+; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(write, argmem: readwrite, inaccessiblemem: none, target_mem: none)
+define internal spir_func void @__device_side_enqueue_block_invoke_2(ptr addrspace(4) noundef readonly captures(none) %.block_descriptor) #6 {
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
+ %block.capture.addr = getelementptr inbounds nuw i8, ptr addrspace(4) %.block_descriptor, i64 28
+ %0 = load i8, ptr addrspace(4) %block.capture.addr, align 4, !tbaa !13
%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
+ %block.capture.addr1 = getelementptr inbounds nuw i8, ptr addrspace(4) %.block_descriptor, i64 16
+ %1 = load ptr addrspace(1), ptr addrspace(4) %block.capture.addr1, align 8, !tbaa !10
+ %block.capture.addr2 = getelementptr inbounds nuw i8, ptr addrspace(4) %.block_descriptor, i64 24
+ %2 = load i32, ptr addrspace(4) %block.capture.addr2, align 8, !tbaa !2
+ %idxprom = sext i32 %2 to i64
+ %arrayidx = getelementptr inbounds [4 x i8], ptr addrspace(1) %1, i64 %idxprom
+ store i32 %conv, ptr addrspace(1) %arrayidx, align 4, !tbaa !2
ret void
}
-define spir_kernel void @__device_side_enqueue_block_invoke_kernel(ptr addrspace(4) %0) {
+; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(write, argmem: readwrite, inaccessiblemem: none, target_mem: none)
+define internal spir_kernel void @__device_side_enqueue_block_invoke_2_kernel(ptr addrspace(4) readonly captures(none) %0) #6 {
entry:
- call spir_func void @__device_side_enqueue_block_invoke(ptr addrspace(4) %0)
+ %block.capture.addr.i = getelementptr inbounds nuw i8, ptr addrspace(4) %0, i64 28
+ %1 = load i8, ptr addrspace(4) %block.capture.addr.i, align 4, !tbaa !13
+ %conv.i = sext i8 %1 to i32
+ %block.capture.addr1.i = getelementptr inbounds nuw i8, ptr addrspace(4) %0, i64 16
+ %2 = load ptr addrspace(1), ptr addrspace(4) %block.capture.addr1.i, align 8, !tbaa !10
+ %block.capture.addr2.i = getelementptr inbounds nuw i8, ptr addrspace(4) %0, i64 24
+ %3 = load i32, ptr addrspace(4) %block.capture.addr2.i, align 8, !tbaa !2
+ %idxprom.i = sext i32 %3 to i64
+ %arrayidx.i = getelementptr inbounds [4 x i8], ptr addrspace(1) %2, i64 %idxprom.i
+ store i32 %conv.i, ptr addrspace(1) %arrayidx.i, align 4, !tbaa !2
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(target("spirv.Queue"), i32, ptr, ptr addrspace(4), ptr addrspace(4)) local_unnamed_addr
-define internal spir_func void @__device_side_enqueue_block_invoke_2(ptr addrspace(4) noundef %.block_descriptor) {
+; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(readwrite, inaccessiblemem: none, target_mem: none)
+define internal spir_func void @__device_side_enqueue_block_invoke_3(ptr addrspace(4) noundef readonly captures(none) %.block_descriptor) #7 {
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
+ %block.capture.addr = getelementptr inbounds nuw i8, ptr addrspace(4) %.block_descriptor, i64 24
+ %0 = load ptr addrspace(1), ptr addrspace(4) %block.capture.addr, align 8, !tbaa !10
+ %block.capture.addr1 = getelementptr inbounds nuw i8, ptr addrspace(4) %.block_descriptor, i64 32
+ %1 = load i32, ptr addrspace(4) %block.capture.addr1, align 8, !tbaa !2
+ %idxprom = sext i32 %1 to i64
+ %arrayidx = getelementptr inbounds [4 x i8], ptr addrspace(1) %0, i64 %idxprom
+ %2 = load i32, ptr addrspace(1) %arrayidx, align 4, !tbaa !2
+ %block.capture.addr2 = getelementptr inbounds nuw i8, ptr addrspace(4) %.block_descriptor, i64 16
+ %3 = load ptr addrspace(1), ptr addrspace(4) %block.capture.addr2, align 8, !tbaa !10
+ %arrayidx5 = getelementptr inbounds [4 x i8], ptr addrspace(1) %3, i64 %idxprom
+ store i32 %2, ptr addrspace(1) %arrayidx5, align 4, !tbaa !2
ret void
}
-define spir_kernel void @__device_side_enqueue_block_invoke_2_kernel(ptr addrspace(4) %0) {
+; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(readwrite, inaccessiblemem: none, target_mem: none)
+define internal spir_kernel void @__device_side_enqueue_block_invoke_3_kernel(ptr addrspace(4) readonly captures(none) %0) #7 {
entry:
- call spir_func void @__device_side_enqueue_block_invoke_2(ptr addrspace(4) %0)
+ %block.capture.addr.i = getelementptr inbounds nuw i8, ptr addrspace(4) %0, i64 24
+ %1 = load ptr addrspace(1), ptr addrspace(4) %block.capture.addr.i, align 8, !tbaa !10
+ %block.capture.addr1.i = getelementptr inbounds nuw i8, ptr addrspace(4) %0, i64 32
+ %2 = load i32, ptr addrspace(4) %block.capture.addr1.i, align 8, !tbaa !2
+ %idxprom.i = sext i32 %2 to i64
+ %arrayidx.i = getelementptr inbounds [4 x i8], ptr addrspace(1) %1, i64 %idxprom.i
+ %3 = load i32, ptr addrspace(1) %arrayidx.i, align 4, !tbaa !2
+ %block.capture.addr2.i = getelementptr inbounds nuw i8, ptr addrspace(4) %0, i64 16
+ %4 = load ptr addrspace(1), ptr addrspace(4) %block.capture.addr2.i, align 8, !tbaa !10
+ %arrayidx5.i = getelementptr inbounds [4 x i8], ptr addrspace(1) %4, i64 %idxprom.i
+ store i32 %3, ptr addrspace(1) %arrayidx5.i, align 4, !tbaa !2
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))
-
-define internal spir_func void @__device_side_enqueue_block_invoke_3(ptr addrspace(4) noundef %.block_descriptor, ptr addrspace(3) noundef %p) {
+; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(none)
+define internal spir_func void @__device_side_enqueue_block_invoke_4(ptr addrspace(4) readnone captures(none) %.block_descriptor, ptr addrspace(3) readnone captures(none) %p) #5 {
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) {
+; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(none)
+define internal spir_kernel void @__device_side_enqueue_block_invoke_4_kernel(ptr addrspace(4) readnone captures(none) %0, ptr addrspace(3) readnone captures(none) %1) #5 {
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) local_unnamed_addr
+
+; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite)
+declare void @llvm.lifetime.end.p0(ptr captures(none)) #2
-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) {
+; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(none)
+define internal spir_func void @__device_side_enqueue_block_invoke_5(ptr addrspace(4) readnone captures(none) %.block_descriptor, ptr addrspace(3) readnone captures(none) %p1, ptr addrspace(3) readnone captures(none) %p2, ptr addrspace(3) readnone captures(none) %p3) #5 {
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) {
+; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(none)
+define internal spir_kernel void @__device_side_enqueue_block_invoke_5_kernel(ptr addrspace(4) readnone captures(none) %0, ptr addrspace(3) readnone captures(none) %1, ptr addrspace(3) readnone captures(none) %2, ptr addrspace(3) readnone captures(none) %3) #5 {
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) local_unnamed_addr
-define internal spir_func void @__device_side_enqueue_block_invoke_5(ptr addrspace(4) noundef %.block_descriptor) {
+; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(readwrite, inaccessiblemem: none, target_mem: none)
+define internal spir_func void @__device_side_enqueue_block_invoke_6(ptr addrspace(4) noundef readonly captures(none) %.block_descriptor) #7 {
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
+ %block.capture.addr = getelementptr inbounds nuw i8, ptr addrspace(4) %.block_descriptor, i64 24
+ %0 = load ptr addrspace(1), ptr addrspace(4) %block.capture.addr, align 8, !tbaa !10
+ %block.capture.addr1 = getelementptr inbounds nuw i8, ptr addrspace(4) %.block_descriptor, i64 32
+ %1 = load i32, ptr addrspace(4) %block.capture.addr1, align 8, !tbaa !2
+ %idxprom = sext i32 %1 to i64
+ %arrayidx = getelementptr inbounds [4 x i8], ptr addrspace(1) %0, i64 %idxprom
+ %2 = load i32, ptr addrspace(1) %arrayidx, align 4, !tbaa !2
+ %block.capture.addr2 = getelementptr inbounds nuw i8, ptr addrspace(4) %.block_descriptor, i64 16
+ %3 = load ptr addrspace(1), ptr addrspace(4) %block.capture.addr2, align 8, !tbaa !10
+ %arrayidx5 = getelementptr inbounds [4 x i8], ptr addrspace(1) %3, i64 %idxprom
+ store i32 %2, ptr addrspace(1) %arrayidx5, align 4, !tbaa !2
ret void
}
-define spir_kernel void @__device_side_enqueue_block_invoke_5_kernel(ptr addrspace(4) %0) {
+; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(readwrite, inaccessiblemem: none, target_mem: none)
+define internal spir_kernel void @__device_side_enqueue_block_invoke_6_kernel(ptr addrspace(4) readonly captures(none) %0) #7 {
entry:
- call spir_func void @__device_side_enqueue_block_invoke_5(ptr addrspace(4) %0)
+ %block.capture.addr.i = getelementptr inbounds nuw i8, ptr addrspace(4) %0, i64 24
+ %1 = load ptr addrspace(1), ptr addrspace(4) %block.capture.addr.i, align 8, !tbaa !10
+ %block.capture.addr1.i = getelementptr inbounds nuw i8, ptr addrspace(4) %0, i64 32
+ %2 = load i32, ptr addrspace(4) %block.capture.addr1.i, align 8, !tbaa !2
+ %idxprom.i = sext i32 %2 to i64
+ %arrayidx.i = getelementptr inbounds [4 x i8], ptr addrspace(1) %1, i64 %idxprom.i
+ %3 = load i32, ptr addrspace(1) %arrayidx.i, align 4, !tbaa !2
+ %block.capture.addr2.i = getelementptr inbounds nuw i8, ptr addrspace(4) %0, i64 16
+ %4 = load ptr addrspace(1), ptr addrspace(4) %block.capture.addr2.i, align 8, !tbaa !10
+ %arrayidx5.i = getelementptr inbounds [4 x i8], ptr addrspace(1) %4, i64 %idxprom.i
+ store i32 %3, ptr addrspace(1) %arrayidx5.i, align 4, !tbaa !2
ret void
}
+
+attributes #0 = { "objc_arc_inert" }
+attributes #1 = { convergent norecurse nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #2 = { nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }
+attributes #3 = { nocallback nofree nounwind willreturn memory(argmem: readwrite) }
+attributes #4 = { convergent nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #5 = { mustprogress nofree norecurse nosync nounwind willreturn memory(none) "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #6 = { mustprogress nofree norecurse nosync nounwind willreturn memory(write, argmem: readwrite, inaccessiblemem: none, target_mem: none) "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #7 = { mustprogress nofree norecurse nosync nounwind willreturn memory(readwrite, inaccessiblemem: none, target_mem: none) "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #8 = { nounwind }
+attributes #9 = { convergent nounwind }
+
+!opencl.ocl.version = !{!0}
+!llvm.errno.tbaa = !{!2}
+
+!0 = !{i32 3, i32 0}
+!2 = !{!3, !3, i64 0}
+!3 = !{!"int", !4, i64 0}
+!4 = !{!"omnipotent char", !5, i64 0}
+!5 = !{!"Simple C/C++ TBAA"}
+!6 = !{i32 1, i32 1, i32 0, i32 0}
+!7 = !{!"none", !"none", !"none", !"none"}
+!8 = !{!"int*", !"int*", !"int", !"char"}
+!9 = !{!"", !"", !"", !""}
+!10 = !{!11, !11, i64 0}
+!11 = !{!"p1 int", !12, i64 0}
+!12 = !{!"any pointer", !4, i64 0}
+!13 = !{!4, !4, i64 0}
>From 8d8a3d641adcc18d151a792b37b07d888f9d077b Mon Sep 17 00:00:00 2001
From: idubinov <igor.dubinov at amd.com>
Date: Tue, 7 Apr 2026 11:03:51 -0500
Subject: [PATCH 12/12] undo const cast
---
llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp | 7 +++----
1 file changed, 3 insertions(+), 4 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
index 95f4466ad94e3..119b1ae63a2de 100644
--- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
@@ -2803,9 +2803,8 @@ static bool buildEnqueueKernel(const SPIRV::IncomingCall *Call,
const SPIRVTypeInst Int8Ty = GR->getOrCreateSPIRVIntegerType(8, MIRBuilder);
const SPIRVTypeInst Int8PtrGen = GR->getOrCreateSPIRVPointerType(
Int8Ty, MIRBuilder, SPIRV::StorageClass::Generic);
- const Type *PType = getBlockStructType(BlockLiteralReg, MRI);
+ Type *PType = const_cast<Type *>(getBlockStructType(BlockLiteralReg, MRI));
- /// TODO: check if already Int8Ptr
Register ParamReg =
createVirtualRegister(Int8PtrGen, GR, MIRBuilder);
MIRBuilder.buildInstr(SPIRV::OpBitcast)
@@ -2813,8 +2812,8 @@ static bool buildEnqueueKernel(const SPIRV::IncomingCall *Call,
.addUse(GR->getSPIRVTypeID(Int8PtrGen))
.addUse(BlockLiteralReg);
/// TODO: these numbers should be obtained from block literal structure.
- Register ParamSizeReg = buildConstantIntReg32(DL.getTypeStoreSize(const_cast<Type *>(PType)), MIRBuilder, GR);
- Register ParamAlignReg = buildConstantIntReg32(DL.getPrefTypeAlign(const_cast<Type *>(PType)).value(), MIRBuilder, GR);
+ Register ParamSizeReg = buildConstantIntReg32(DL.getTypeStoreSize(PType), MIRBuilder, GR);
+ Register ParamAlignReg = buildConstantIntReg32(DL.getPrefTypeAlign(PType).value(), MIRBuilder, GR);
/// 2.4 Local Size Array
More information about the llvm-commits
mailing list