[llvm] [SPIRV] Add support for SPV_KHR_abort extension (PR #193037)
Victor Mustya via llvm-commits
llvm-commits at lists.llvm.org
Tue May 5 07:23:55 PDT 2026
https://github.com/vmustya updated https://github.com/llvm/llvm-project/pull/193037
>From eab1954fec17b572496375bcb09bff55a5e68782 Mon Sep 17 00:00:00 2001
From: Victor Mustya <victor.mustya at intel.com>
Date: Mon, 20 Apr 2026 10:23:55 -0700
Subject: [PATCH 1/7] [SPIRV] Add support for SPV_KHR_abort extension
This commit adds support for the SPV_KHR_abort extension in the SPIRV
backend. The extension allows shaders to abort execution with a custom
message.
Assisted-by: Claude Opus 4.7 <noreply at anthropic.com>
---
llvm/docs/SPIRVUsage.rst | 2 +
llvm/include/llvm/IR/IntrinsicsSPIRV.td | 1 +
llvm/lib/Target/SPIRV/SPIRVCommandLine.cpp | 3 +-
llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp | 14 ++++
llvm/lib/Target/SPIRV/SPIRVInstrInfo.td | 2 +
.../Target/SPIRV/SPIRVInstructionSelector.cpp | 61 ++++++++++++++
llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp | 8 ++
.../Target/SPIRV/SPIRVPrepareFunctions.cpp | 60 ++++++++++++++
.../lib/Target/SPIRV/SPIRVSymbolicOperands.td | 5 ++
.../abort-composite-construct.ll | 30 +++++++
.../SPV_KHR_abort/abort-composite.ll | 51 ++++++++++++
.../abort-invalid-message-type.ll | 61 ++++++++++++++
.../SPV_KHR_abort/abort-opencl-cfg.ll | 40 +++++++++
.../SPV_KHR_abort/abort-opencl-source.ll | 81 +++++++++++++++++++
.../extensions/SPV_KHR_abort/abort-opencl.ll | 62 ++++++++++++++
.../SPIRV/extensions/SPV_KHR_abort/abort.ll | 24 ++++++
16 files changed, 504 insertions(+), 1 deletion(-)
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-composite-construct.ll
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-composite.ll
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-invalid-message-type.ll
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-opencl-cfg.ll
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-opencl-source.ll
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-opencl.ll
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort.ll
diff --git a/llvm/docs/SPIRVUsage.rst b/llvm/docs/SPIRVUsage.rst
index 75a477b86da05..e66f8bb374a94 100644
--- a/llvm/docs/SPIRVUsage.rst
+++ b/llvm/docs/SPIRVUsage.rst
@@ -213,6 +213,8 @@ Below is a list of supported SPIR-V extensions, sorted alphabetically by their e
- Allows to allocate local arrays whose number of elements is unknown at compile time.
* - ``SPV_INTEL_joint_matrix``
- Adds few matrix capabilities on top of SPV_KHR_cooperative_matrix extension, such as matrix prefetch, get element coordinate and checked load/store/construct instructions, tensor float 32 and bfloat type interpretations for multiply-add instruction.
+ * - ``SPV_KHR_abort``
+ - Adds a function-termination instruction that signals an abnormal error to the client API. Currently supported in OpenCL environments only.
* - ``SPV_KHR_bit_instructions``
- Enables bit instructions to be used by SPIR-V modules without requiring the Shader capability.
* - ``SPV_KHR_expect_assume``
diff --git a/llvm/include/llvm/IR/IntrinsicsSPIRV.td b/llvm/include/llvm/IR/IntrinsicsSPIRV.td
index d2a5fa1f08724..c4733d3942f2e 100644
--- a/llvm/include/llvm/IR/IntrinsicsSPIRV.td
+++ b/llvm/include/llvm/IR/IntrinsicsSPIRV.td
@@ -39,6 +39,7 @@ let TargetPrefix = "spv" in {
def int_spv_selection_merge : Intrinsic<[], [llvm_any_ty, llvm_i32_ty], [ImmArg<ArgIndex<1>>]>;
def int_spv_cmpxchg : Intrinsic<[llvm_i32_ty], [llvm_any_ty, llvm_vararg_ty]>;
def int_spv_unreachable : Intrinsic<[], []>;
+ def int_spv_abort : Intrinsic<[], [llvm_vararg_ty], [IntrNoReturn]>;
def int_spv_alloca : Intrinsic<[llvm_any_ty], [llvm_i32_ty], [ImmArg<ArgIndex<0>>]>;
def int_spv_alloca_array : Intrinsic<[llvm_any_ty], [llvm_anyint_ty, llvm_i32_ty], [ImmArg<ArgIndex<1>>]>;
def int_spv_undef : Intrinsic<[llvm_i32_ty], []>;
diff --git a/llvm/lib/Target/SPIRV/SPIRVCommandLine.cpp b/llvm/lib/Target/SPIRV/SPIRVCommandLine.cpp
index 734a03ff60141..863c3b6d1f757 100644
--- a/llvm/lib/Target/SPIRV/SPIRVCommandLine.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVCommandLine.cpp
@@ -176,7 +176,8 @@ static const StringMap<SPIRV::Extension::Extension> SPIRVExtensionMap = {
{"SPV_EXT_image_raw10_raw12",
SPIRV::Extension::Extension::SPV_EXT_image_raw10_raw12},
{"SPV_INTEL_unstructured_loop_controls",
- SPIRV::Extension::Extension::SPV_INTEL_unstructured_loop_controls}};
+ SPIRV::Extension::Extension::SPV_INTEL_unstructured_loop_controls},
+ {"SPV_KHR_abort", SPIRV::Extension::Extension::SPV_KHR_abort}};
bool SPIRVExtensionsParser::parse(cl::Option &O, StringRef ArgName,
StringRef ArgValue, ExtensionSet &Vals) {
diff --git a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
index 7a886de005b88..6e109c1dd74ab 100644
--- a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
@@ -2253,6 +2253,20 @@ Instruction *SPIRVEmitIntrinsics::visitAtomicCmpXchgInst(AtomicCmpXchgInst &I) {
Instruction *SPIRVEmitIntrinsics::visitUnreachableInst(UnreachableInst &I) {
IRBuilder<> B(I.getParent());
B.SetInsertPoint(&I);
+ // OpAbortKHR is itself a SPIR-V block terminator. If the previous instruction
+ // is a call to llvm.spv.abort, do not emit an additional OpUnreachable, which
+ // would leave the block with two terminators and produce invalid SPIR-V.
+ for (Instruction *Prev = I.getPrevNode(); Prev; Prev = Prev->getPrevNode()) {
+ if (Prev->isDebugOrPseudoInst())
+ continue;
+ auto *CI = dyn_cast<CallInst>(Prev);
+ if (!CI)
+ break;
+ Intrinsic::ID IID = CI->getIntrinsicID();
+ if (IID == Intrinsic::spv_abort)
+ return &I;
+ break;
+ }
B.CreateIntrinsic(Intrinsic::spv_unreachable, {});
return &I;
}
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstrInfo.td b/llvm/lib/Target/SPIRV/SPIRVInstrInfo.td
index 04ff442bbeed2..a1162f2fbc664 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstrInfo.td
+++ b/llvm/lib/Target/SPIRV/SPIRVInstrInfo.td
@@ -641,6 +641,8 @@ let isReturn = 1, hasDelaySlot = 0, isBarrier = 0, isTerminator = 1, isNotDuplic
def OpReturn: SimpleOp<"OpReturn", 253>;
def OpReturnValue: Op<254, (outs), (ins ID:$ret), "OpReturnValue $ret">;
def OpUnreachable: SimpleOp<"OpUnreachable", 255>;
+ def OpAbortKHR: Op<5121, (outs), (ins TYPE:$msg_type, ID:$msg),
+ "OpAbortKHR $msg_type $msg">;
}
def OpLifetimeStart: Op<256, (outs), (ins ID:$ptr, i32imm:$sz), "OpLifetimeStart $ptr $sz">;
def OpLifetimeStop: Op<257, (outs), (ins ID:$ptr, i32imm:$sz), "OpLifetimeStop $ptr $sz">;
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index f1e0450bb20f9..7414880ecb05b 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -350,6 +350,7 @@ class SPIRVInstructionSelector : public InstructionSelector {
bool diagnoseUnsupported(const MachineInstr &I, const Twine &Msg) const;
+ bool selectAbort(MachineInstr &I) const;
bool selectFrameIndex(Register ResVReg, SPIRVTypeInst ResType,
MachineInstr &I) const;
bool selectAllocaArray(Register ResVReg, SPIRVTypeInst ResType,
@@ -4635,6 +4636,8 @@ bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpUnreachable))
.constrainAllUses(TII, TRI, RBI);
return true;
+ case Intrinsic::spv_abort:
+ return selectAbort(I);
case Intrinsic::spv_alloca:
return selectFrameIndex(ResVReg, ResType, I);
case Intrinsic::spv_alloca_array:
@@ -6250,6 +6253,64 @@ bool SPIRVInstructionSelector::selectAllocaArray(Register ResVReg,
return true;
}
+bool SPIRVInstructionSelector::selectAbort(MachineInstr &I) const {
+ if (!STI.canUseExtension(SPIRV::Extension::SPV_KHR_abort))
+ report_fatal_error("OpAbortKHR instruction requires the following "
+ "SPIR-V extension: SPV_KHR_abort",
+ false);
+ // The intrinsic is declared as variadic so it can carry composite message
+ // types (vectors and structs) without LLVM IR mangling restrictions, but
+ // OpAbortKHR takes exactly one Message operand.
+ if (I.getNumExplicitOperands() != 2)
+ report_fatal_error("llvm.spv.abort must be called with exactly one "
+ "message argument",
+ false);
+ Register MsgReg = I.getOperand(1).getReg();
+ SPIRVTypeInst MsgType = GR.getSPIRVTypeForVReg(MsgReg);
+ assert(MsgType && "Message argument of llvm.spv.abort has no SPIR-V type");
+ // SPV_KHR_abort requires Message Type to be a concrete type. Per the
+ // SPIR-V "Concrete Type" definition, that means a numerical scalar
+ // (int/float), a vector, matrix, or any aggregate (array/struct)
+ // recursively containing only such types. OpTypeBool, OpTypeVoid,
+ // pointers, opaque handles and similar abstract/non-concrete types are
+ // rejected up front rather than emitting invalid SPIR-V. Validate
+ // recursively so that e.g. a struct containing a bool or pointer is also
+ // rejected.
+ SmallVector<SPIRVTypeInst, 4> Worklist{MsgType};
+ while (!Worklist.empty()) {
+ SPIRVTypeInst Ty = Worklist.pop_back_val();
+ switch (Ty->getOpcode()) {
+ case SPIRV::OpTypeInt:
+ case SPIRV::OpTypeFloat:
+ break;
+ case SPIRV::OpTypeVector:
+ case SPIRV::OpTypeMatrix:
+ case SPIRV::OpTypeArray:
+ // Operand 1 holds the element/component type id.
+ Worklist.push_back(
+ GR.getSPIRVTypeForVReg(Ty->getOperand(1).getReg()));
+ break;
+ case SPIRV::OpTypeStruct:
+ // Operands 1..N hold the field type ids.
+ for (unsigned Idx = 1, E = Ty->getNumOperands(); Idx < E; ++Idx)
+ Worklist.push_back(
+ GR.getSPIRVTypeForVReg(Ty->getOperand(Idx).getReg()));
+ break;
+ default:
+ report_fatal_error("llvm.spv.abort message type must be a concrete "
+ "SPIR-V type (numerical scalar, vector, matrix, "
+ "or aggregate of such types)",
+ false);
+ }
+ }
+ MachineBasicBlock &BB = *I.getParent();
+ BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpAbortKHR))
+ .addUse(GR.getSPIRVTypeID(MsgType))
+ .addUse(MsgReg)
+ .constrainAllUses(TII, TRI, RBI);
+ return true;
+}
+
bool SPIRVInstructionSelector::selectFrameIndex(Register ResVReg,
SPIRVTypeInst ResType,
MachineInstr &I) const {
diff --git a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
index aec9b15df9189..5412e50d6afc5 100644
--- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
@@ -1945,6 +1945,14 @@ void addInstrRequirements(const MachineInstr &MI,
Reqs.addExtension(SPIRV::Extension::SPV_KHR_shader_clock);
Reqs.addCapability(SPIRV::Capability::ShaderClockKHR);
break;
+ case SPIRV::OpAbortKHR:
+ if (!ST.canUseExtension(SPIRV::Extension::SPV_KHR_abort))
+ report_fatal_error("OpAbortKHR instruction requires the "
+ "following SPIR-V extension: SPV_KHR_abort",
+ false);
+ Reqs.addExtension(SPIRV::Extension::SPV_KHR_abort);
+ Reqs.addCapability(SPIRV::Capability::AbortKHR);
+ break;
case SPIRV::OpFunctionPointerCallINTEL:
if (ST.canUseExtension(SPIRV::Extension::SPV_INTEL_function_pointers)) {
Reqs.addExtension(SPIRV::Extension::SPV_INTEL_function_pointers);
diff --git a/llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp b/llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp
index a3b44ad6d31d5..7e955e2d6d2df 100644
--- a/llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp
@@ -19,6 +19,7 @@
//===----------------------------------------------------------------------===//
#include "SPIRV.h"
+#include "SPIRVBuiltins.h"
#include "SPIRVSubtarget.h"
#include "SPIRVTargetMachine.h"
#include "SPIRVUtils.h"
@@ -33,6 +34,7 @@
#include "llvm/IR/Intrinsics.h"
#include "llvm/IR/IntrinsicsSPIRV.h"
#include "llvm/Transforms/Utils/Cloning.h"
+#include "llvm/Transforms/Utils/Local.h"
#include "llvm/Transforms/Utils/LowerMemIntrinsics.h"
#include <regex>
@@ -43,6 +45,7 @@ namespace {
class SPIRVPrepareFunctions : public ModulePass {
const SPIRVTargetMachine &TM;
bool substituteIntrinsicCalls(Function *F);
+ bool substituteAbortKHRCalls(Function *F);
Function *removeAggregateTypesFromSignature(Function *F);
bool removeAggregateTypesFromCalls(Function *F);
@@ -595,6 +598,62 @@ SPIRVPrepareFunctions::removeAggregateTypesFromSignature(Function *F) {
return NewF;
}
+// Replace OpenCL/SPIR-V style calls to `__spirv_AbortKHR(message)` with calls
+// to the `llvm.spv.abort` target intrinsic, so that they go through the same
+// instruction-selection path as the intrinsic and get lowered to OpAbortKHR.
+bool SPIRVPrepareFunctions::substituteAbortKHRCalls(Function *F) {
+ if (F->isDeclaration())
+ return false;
+
+ SmallVector<CallInst *> Calls;
+ for (Instruction &I : instructions(F)) {
+ auto *CI = dyn_cast<CallInst>(&I);
+ if (!CI)
+ continue;
+ Function *Callee = CI->getCalledFunction();
+ if (!Callee || Callee->isIntrinsic())
+ continue;
+ StringRef Demangled = Callee->getName();
+ std::string DemangledStr = getOclOrSpirvBuiltinDemangledName(Demangled);
+ if (DemangledStr.empty())
+ continue;
+ std::string BuiltinName = SPIRV::lookupBuiltinNameHelper(DemangledStr);
+ if (StringRef(BuiltinName) != "__spirv_AbortKHR")
+ continue;
+ if (CI->arg_size() != 1)
+ continue;
+ Calls.push_back(CI);
+ }
+
+ if (Calls.empty())
+ return false;
+
+ for (CallInst *CI : Calls) {
+ IRBuilder<> B(CI);
+ Value *Msg = CI->getArgOperand(0);
+ // The OpenCL C ABI may pass aggregate arguments by pointer (byval). In
+ // that case load the underlying value so that OpAbortKHR receives the
+ // composite itself, as required by the SPV_KHR_abort spec ("Message Type
+ // must be a concrete type").
+ if (CI->isByValArgument(0)) {
+ Type *AggTy = CI->getParamByValType(0);
+ Msg = B.CreateLoad(AggTy, Msg);
+ }
+ B.CreateIntrinsic(Intrinsic::spv_abort, {}, {Msg});
+ // OpAbortKHR is itself a SPIR-V function-termination instruction and must
+ // be the last instruction in its block. Drop the original call and
+ // everything that follows it (the OpenCL ABI typically appends stores into
+ // the return slot and a `ret`), and re-terminate the block with
+ // `unreachable`. We use changeToUnreachable so that any successor PHI
+ // nodes have the now-removed predecessor edge cleaned up; otherwise the
+ // IR verifier would reject mismatched PHI incoming entries. The matching
+ // suppression in SPIRVEmitIntrinsics::visitUnreachableInst ensures no
+ // extra OpUnreachable is emitted after OpAbortKHR.
+ changeToUnreachable(CI);
+ }
+ return true;
+}
+
// Mutates indirect callsites iff if aggregate argument/return types are present
// with the types replaced by i32 types. The change in types is noted in
// 'spv.mutated_callsites' metadata for later restoration.
@@ -684,6 +743,7 @@ bool SPIRVPrepareFunctions::runOnModule(Module &M) {
for (Function &F : M) {
Changed |= substituteIntrinsicCalls(&F);
+ Changed |= substituteAbortKHRCalls(&F);
Changed |= sortBlocks(F);
Changed |= removeAggregateTypesFromCalls(&F);
}
diff --git a/llvm/lib/Target/SPIRV/SPIRVSymbolicOperands.td b/llvm/lib/Target/SPIRV/SPIRVSymbolicOperands.td
index 56608a80f4b23..68ee90f2c6792 100644
--- a/llvm/lib/Target/SPIRV/SPIRVSymbolicOperands.td
+++ b/llvm/lib/Target/SPIRV/SPIRVSymbolicOperands.td
@@ -398,6 +398,10 @@ defm SPV_EXT_image_raw10_raw12 :ExtensionOperand<133, [EnvOpenCL, EnvVulkan]>;
defm SPV_ALTERA_arbitrary_precision_floating_point: ExtensionOperand<134, [EnvOpenCL]>;
defm SPV_KHR_fma : ExtensionOperand<135, [EnvVulkan, EnvOpenCL]>;
defm SPV_INTEL_masked_gather_scatter : ExtensionOperand<136, [EnvOpenCL]>;
+// SPV_KHR_abort is currently only supported in the OpenCL environment.
+// Vulkan/Shader support requires emitting explicit layout decorations on the
+// Message Type, which is not yet implemented.
+defm SPV_KHR_abort : ExtensionOperand<137, [EnvOpenCL]>;
//===----------------------------------------------------------------------===//
// Multiclass used to define Capabilities enum values and at the same time
@@ -598,6 +602,7 @@ defm GlobalVariableFPGADecorationsINTEL : CapabilityOperand<6189, 0, 0, [SPV_INT
defm CacheControlsINTEL : CapabilityOperand<6441, 0, 0, [SPV_INTEL_cache_controls], []>;
defm CooperativeMatrixKHR : CapabilityOperand<6022, 0, 0, [SPV_KHR_cooperative_matrix], []>;
defm ArithmeticFenceEXT : CapabilityOperand<6144, 0, 0, [SPV_EXT_arithmetic_fence], []>;
+defm AbortKHR : CapabilityOperand<5120, 0, 0, [SPV_KHR_abort], []>;
defm SplitBarrierINTEL : CapabilityOperand<6141, 0, 0, [SPV_INTEL_split_barrier], []>;
defm CooperativeMatrixCheckedInstructionsINTEL : CapabilityOperand<6192, 0, 0, [SPV_INTEL_joint_matrix], []>;
defm CooperativeMatrixPrefetchINTEL : CapabilityOperand<6411, 0, 0, [SPV_INTEL_joint_matrix], []>;
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-composite-construct.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-composite-construct.ll
new file mode 100644
index 0000000000000..b7ab99b47b0cd
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-composite-construct.ll
@@ -0,0 +1,30 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - -filetype=obj | spirv-val %}
+
+;; Verify that OpAbortKHR can consume a composite produced by
+;; OpCompositeConstruct, mirroring the SPV_KHR_abort spec example which builds
+;; the message via OpCompositeConstruct.
+
+; CHECK-DAG: OpCapability AbortKHR
+; CHECK-DAG: OpExtension "SPV_KHR_abort"
+; CHECK-DAG: %[[#I32:]] = OpTypeInt 32 0
+; CHECK-DAG: %[[#V2:]] = OpTypeVector %[[#I32]] 2
+
+; CHECK: %[[#CC:]] = OpCompositeConstruct %[[#V2]] %{{[0-9]+}} %{{[0-9]+}}
+; CHECK: OpAbortKHR %[[#V2]] %[[#CC]]
+; CHECK-NOT: OpUnreachable
+
+declare void @llvm.spv.abort(...) #0
+
+define spir_kernel void @abort_composite_construct(i32 %a, i32 %b) {
+entry:
+ %va = insertelement <1 x i32> poison, i32 %a, i32 0
+ %vb = insertelement <1 x i32> poison, i32 %b, i32 0
+ %v = shufflevector <1 x i32> %va, <1 x i32> %vb,
+ <2 x i32> <i32 0, i32 1>
+ call void (...) @llvm.spv.abort(<2 x i32> %v)
+ unreachable
+}
+
+attributes #0 = { noreturn }
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-composite.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-composite.ll
new file mode 100644
index 0000000000000..f8c779dc7b891
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-composite.ll
@@ -0,0 +1,51 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - -filetype=obj | spirv-val %}
+
+; Verify that OpAbortKHR can consume a composite (struct or vector) message,
+; built up via OpCompositeInsert, mirroring the example in the SPV_KHR_abort
+; specification.
+
+; CHECK-DAG: OpCapability AbortKHR
+; CHECK-DAG: OpExtension "SPV_KHR_abort"
+
+; CHECK-DAG: %[[#I32:]] = OpTypeInt 32 0
+; CHECK-DAG: %[[#F32:]] = OpTypeFloat 32
+; CHECK-DAG: %[[#STRUCT:]] = OpTypeStruct %[[#I32]] %[[#I32]] %[[#F32]]
+; CHECK-DAG: %[[#VEC:]] = OpTypeVector %[[#I32]] 4
+
+; CHECK: %[[#S0:]] = OpCompositeInsert %[[#STRUCT]] %{{[0-9]+}} %{{[0-9]+}} 0
+; CHECK: %[[#S1:]] = OpCompositeInsert %[[#STRUCT]] %{{[0-9]+}} %[[#S0]] 1
+; CHECK: %[[#S2:]] = OpCompositeInsert %[[#STRUCT]] %{{[0-9]+}} %[[#S1]] 2
+; CHECK: OpAbortKHR %[[#STRUCT]] %[[#S2]]
+
+; CHECK: %[[#V0:]] = OpCompositeInsert %[[#VEC]] %{{[0-9]+}} %{{[0-9]+}} 0
+; CHECK: %[[#V1:]] = OpCompositeInsert %[[#VEC]] %{{[0-9]+}} %[[#V0]] 1
+; CHECK: %[[#V2:]] = OpCompositeInsert %[[#VEC]] %{{[0-9]+}} %[[#V1]] 2
+; CHECK: %[[#V3:]] = OpCompositeInsert %[[#VEC]] %{{[0-9]+}} %[[#V2]] 3
+; CHECK: OpAbortKHR %[[#VEC]] %[[#V3]]
+
+%struct.Msg = type { i32, i32, float }
+
+declare void @llvm.spv.abort(...) #0
+
+define spir_kernel void @abort_with_struct(i32 %x, i32 %y, float %z) {
+entry:
+ %m0 = insertvalue %struct.Msg poison, i32 %x, 0
+ %m1 = insertvalue %struct.Msg %m0, i32 %y, 1
+ %m2 = insertvalue %struct.Msg %m1, float %z, 2
+ call void (...) @llvm.spv.abort(%struct.Msg %m2)
+ unreachable
+}
+
+define spir_kernel void @abort_with_vector(i32 %a, i32 %b, i32 %c, i32 %d) {
+entry:
+ %v0 = insertelement <4 x i32> poison, i32 %a, i32 0
+ %v1 = insertelement <4 x i32> %v0, i32 %b, i32 1
+ %v2 = insertelement <4 x i32> %v1, i32 %c, i32 2
+ %v3 = insertelement <4 x i32> %v2, i32 %d, i32 3
+ call void (...) @llvm.spv.abort(<4 x i32> %v3)
+ unreachable
+}
+
+attributes #0 = { noreturn }
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-invalid-message-type.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-invalid-message-type.ll
new file mode 100644
index 0000000000000..d491ac70d80de
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-invalid-message-type.ll
@@ -0,0 +1,61 @@
+; RUN: split-file %s %t
+; RUN: not llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %t/top-level.ll -o /dev/null 2>&1 | FileCheck %s
+; RUN: not llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %t/nested-struct.ll -o /dev/null 2>&1 | FileCheck %s
+; RUN: not llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %t/nested-array.ll -o /dev/null 2>&1 | FileCheck %s
+; RUN: not llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %t/top-level-bool.ll -o /dev/null 2>&1 | FileCheck %s
+; RUN: not llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %t/nested-bool.ll -o /dev/null 2>&1 | FileCheck %s
+
+; CHECK: llvm.spv.abort message type must be a concrete SPIR-V type
+
+;--- top-level.ll
+declare void @llvm.spv.abort(...) #0
+define void @abort_with_pointer(ptr %p) {
+entry:
+ call void (...) @llvm.spv.abort(ptr %p)
+ unreachable
+}
+attributes #0 = { noreturn }
+
+;--- nested-struct.ll
+%S = type { i32, ptr }
+declare void @llvm.spv.abort(...) #0
+define void @abort_with_struct_of_pointer(ptr %p) {
+entry:
+ %s0 = insertvalue %S poison, i32 0, 0
+ %s1 = insertvalue %S %s0, ptr %p, 1
+ call void (...) @llvm.spv.abort(%S %s1)
+ unreachable
+}
+attributes #0 = { noreturn }
+
+;--- nested-array.ll
+declare void @llvm.spv.abort(...) #0
+define void @abort_with_array_of_pointer(ptr %p) {
+entry:
+ %a0 = insertvalue [2 x ptr] poison, ptr %p, 0
+ %a1 = insertvalue [2 x ptr] %a0, ptr %p, 1
+ call void (...) @llvm.spv.abort([2 x ptr] %a1)
+ unreachable
+}
+attributes #0 = { noreturn }
+
+;--- top-level-bool.ll
+declare void @llvm.spv.abort(...) #0
+define void @abort_with_bool(i1 %b) {
+entry:
+ call void (...) @llvm.spv.abort(i1 %b)
+ unreachable
+}
+attributes #0 = { noreturn }
+
+;--- nested-bool.ll
+%B = type { i32, i1 }
+declare void @llvm.spv.abort(...) #0
+define void @abort_with_struct_of_bool(i1 %b) {
+entry:
+ %s0 = insertvalue %B poison, i32 0, 0
+ %s1 = insertvalue %B %s0, i1 %b, 1
+ call void (...) @llvm.spv.abort(%B %s1)
+ unreachable
+}
+attributes #0 = { noreturn }
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-opencl-cfg.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-opencl-cfg.ll
new file mode 100644
index 0000000000000..47169b3c84e1d
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-opencl-cfg.ll
@@ -0,0 +1,40 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - -filetype=obj | spirv-val %}
+
+;; The OpenCL builtin rewrite must not break IR when the abort call sits in a
+;; conditional block whose successor has PHI nodes referring back to it. The
+;; pass should drop the original successor edge from the abort block (along
+;; with the trailing branch), so the merge block's PHI is left consistent.
+;; Regression test for a verifier failure where a stale PHI predecessor was
+;; left after the abort block was re-terminated with `unreachable`.
+
+; CHECK-DAG: OpCapability AbortKHR
+; CHECK-DAG: OpExtension "SPV_KHR_abort"
+
+; CHECK-DAG: %[[#I32:]] = OpTypeInt 32 0
+
+; The abort block must end with OpAbortKHR (no OpReturn / OpBranch / OpUnreachable
+; in between).
+; CHECK: OpAbortKHR %[[#I32]] %{{[0-9]+}}
+; CHECK-NOT: OpUnreachable
+
+declare spir_func void @_Z16__spirv_AbortKHRj(i32) #0
+
+define spir_kernel void @abort_in_conditional(i1 %c, i32 %x) {
+entry:
+ br i1 %c, label %then, label %else
+
+then:
+ call spir_func void @_Z16__spirv_AbortKHRj(i32 %x)
+ br label %merge
+
+else:
+ br label %merge
+
+merge:
+ %v = phi i32 [ 1, %then ], [ 2, %else ]
+ ret void
+}
+
+attributes #0 = { convergent nounwind }
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-opencl-source.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-opencl-source.ll
new file mode 100644
index 0000000000000..2a4417fa3ace7
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-opencl-source.ll
@@ -0,0 +1,81 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - -filetype=obj | spirv-val %}
+
+;; This test models the actual LLVM IR that Clang produces for OpenCL C calls
+;; to `__spirv_AbortKHR`, including the trailing `ret void` (the OpenCL ABI
+;; emits the call as a regular spir_func call followed by a return). The
+;; backend must drop the trailing return because `OpAbortKHR` is itself a
+;; SPIR-V function-termination instruction; the resulting block must have no
+;; instructions after `OpAbortKHR`.
+;;
+;; Source (compiled with: clang -cc1 -triple spir64-unknown-unknown
+;; -cl-std=CL2.0 -finclude-default-header -emit-llvm -O0):
+;;
+;; void __spirv_AbortKHR(uint);
+;; __kernel void k_scalar(uint x) { __spirv_AbortKHR(x); }
+;;
+;; typedef struct { uint a; uint b; uint c; } Msg;
+;; void __spirv_AbortKHR(Msg);
+;; __kernel void k_struct(uint a, uint b, uint c) {
+;; Msg m = { a, b, c };
+;; __spirv_AbortKHR(m);
+;; }
+
+; CHECK-DAG: OpCapability AbortKHR
+; CHECK-DAG: OpExtension "SPV_KHR_abort"
+
+; CHECK-DAG: %[[#I32:]] = OpTypeInt 32 0
+; CHECK-DAG: %[[#STRUCT:]] = OpTypeStruct %[[#I32]] %[[#I32]] %[[#I32]]
+
+; Scalar argument: passed by value, lowered to OpAbortKHR with the i32 type.
+; The trailing OpReturn from the OpenCL ABI must be dropped.
+; CHECK: OpAbortKHR %[[#I32]] %{{[0-9]+}}
+; CHECK-NEXT: OpFunctionEnd
+
+; Struct argument: passed by `byval` pointer, must be loaded so OpAbortKHR
+; receives the composite by value. No OpReturn / OpUnreachable after it.
+; CHECK: %[[#LOADED:]] = OpLoad %[[#STRUCT]]
+; CHECK: OpAbortKHR %[[#STRUCT]] %[[#LOADED]]
+; CHECK-NEXT: OpFunctionEnd
+
+; CHECK-NOT: OpReturn{{[[:space:]]+}}OpFunctionEnd
+; CHECK-NOT: OpUnreachable
+
+%struct.Msg = type { i32, i32, i32 }
+
+declare spir_func void @_Z16__spirv_AbortKHRj(i32) #0
+declare spir_func void @_Z16__spirv_AbortKHR3Msg(ptr byval(%struct.Msg)) #0
+
+define spir_kernel void @k_scalar(i32 noundef %x) {
+entry:
+ %x.addr = alloca i32, align 4
+ store i32 %x, ptr %x.addr, align 4
+ %0 = load i32, ptr %x.addr, align 4
+ call spir_func void @_Z16__spirv_AbortKHRj(i32 noundef %0)
+ ret void
+}
+
+define spir_kernel void @k_struct(i32 noundef %a, i32 noundef %b, i32 noundef %c) {
+entry:
+ %a.addr = alloca i32, align 4
+ %b.addr = alloca i32, align 4
+ %c.addr = alloca i32, align 4
+ %m = alloca %struct.Msg, align 4
+ store i32 %a, ptr %a.addr, align 4
+ store i32 %b, ptr %b.addr, align 4
+ store i32 %c, ptr %c.addr, align 4
+ %p0 = getelementptr inbounds %struct.Msg, ptr %m, i32 0, i32 0
+ %0 = load i32, ptr %a.addr, align 4
+ store i32 %0, ptr %p0, align 4
+ %p1 = getelementptr inbounds %struct.Msg, ptr %m, i32 0, i32 1
+ %1 = load i32, ptr %b.addr, align 4
+ store i32 %1, ptr %p1, align 4
+ %p2 = getelementptr inbounds %struct.Msg, ptr %m, i32 0, i32 2
+ %2 = load i32, ptr %c.addr, align 4
+ store i32 %2, ptr %p2, align 4
+ call spir_func void @_Z16__spirv_AbortKHR3Msg(ptr noundef byval(%struct.Msg) align 4 %m)
+ ret void
+}
+
+attributes #0 = { convergent nounwind }
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-opencl.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-opencl.ll
new file mode 100644
index 0000000000000..331f7f8a41e98
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-opencl.ll
@@ -0,0 +1,62 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - -filetype=obj | spirv-val %}
+
+; Verify that a (mangled or non-mangled) call to the SPIR-V friendly OpenCL
+; built-in `__spirv_AbortKHR` is lowered to OpAbortKHR.
+
+; CHECK-DAG: OpCapability AbortKHR
+; CHECK-DAG: OpExtension "SPV_KHR_abort"
+
+; CHECK-DAG: %[[#I32:]] = OpTypeInt 32 0
+; CHECK-DAG: %[[#NM:]] = OpConstant %[[#I32]] 7
+; CHECK-DAG: %[[#MN:]] = OpConstant %[[#I32]] 11
+; CHECK-DAG: %[[#STRUCT:]] = OpTypeStruct %[[#I32]] %[[#I32]] %[[#I32]]
+
+; CHECK: OpAbortKHR %[[#I32]] %[[#NM]]
+; CHECK-NOT: OpUnreachable
+; CHECK: OpAbortKHR %[[#I32]] %[[#MN]]
+; CHECK-NOT: OpUnreachable
+
+; The OpenCL C ABI passes aggregate arguments by pointer (byval). Verify that
+; such an argument is loaded so OpAbortKHR receives the composite by value.
+; CHECK: %[[#LOADED:]] = OpLoad %[[#STRUCT]]
+; CHECK: OpAbortKHR %[[#STRUCT]] %[[#LOADED]]
+
+%struct.Msg = type { i32, i32, i32 }
+
+; Non-mangled SPIR-V friendly name (commonly used in OpenCL C/C++ via __spirv_*).
+declare void @__spirv_AbortKHR(i32) #0
+
+; Itanium-mangled OpenCL C name: __spirv_AbortKHR(unsigned int).
+declare void @_Z16__spirv_AbortKHRj(i32) #0
+
+; Itanium-mangled OpenCL C name with a struct argument; passed by pointer.
+declare void @_Z16__spirv_AbortKHR3Msg(ptr byval(%struct.Msg)) #0
+
+define spir_kernel void @kernel_nonmangled() {
+entry:
+ call void @__spirv_AbortKHR(i32 7)
+ unreachable
+}
+
+define spir_kernel void @kernel_mangled() {
+entry:
+ call void @_Z16__spirv_AbortKHRj(i32 11)
+ unreachable
+}
+
+define spir_kernel void @kernel_struct(i32 %x, i32 %y, i32 %z) {
+entry:
+ %m = alloca %struct.Msg, align 4
+ %p0 = getelementptr inbounds %struct.Msg, ptr %m, i32 0, i32 0
+ store i32 %x, ptr %p0, align 4
+ %p1 = getelementptr inbounds %struct.Msg, ptr %m, i32 0, i32 1
+ store i32 %y, ptr %p1, align 4
+ %p2 = getelementptr inbounds %struct.Msg, ptr %m, i32 0, i32 2
+ store i32 %z, ptr %p2, align 4
+ call void @_Z16__spirv_AbortKHR3Msg(ptr byval(%struct.Msg) %m)
+ unreachable
+}
+
+attributes #0 = { noreturn }
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort.ll
new file mode 100644
index 0000000000000..a91cb02a36927
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort.ll
@@ -0,0 +1,24 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - -filetype=obj | spirv-val %}
+
+; RUN: not llc -O0 -mtriple=spirv64-unknown-unknown %s -o /dev/null 2>&1 | FileCheck %s --check-prefix=CHECK-ERROR
+
+; CHECK-ERROR: OpAbortKHR instruction requires the following SPIR-V extension: SPV_KHR_abort
+
+; CHECK-DAG: OpCapability AbortKHR
+; CHECK-DAG: OpExtension "SPV_KHR_abort"
+
+; CHECK-DAG: %[[#I32:]] = OpTypeInt 32 0
+; CHECK-DAG: %[[#MSG:]] = OpConstant %[[#I32]] 42
+
+; CHECK: OpAbortKHR %[[#I32]] %[[#MSG]]
+; CHECK-NOT: OpUnreachable
+
+declare void @llvm.spv.abort(...)
+
+define void @abort_with_int() {
+entry:
+ call void (...) @llvm.spv.abort(i32 42)
+ unreachable
+}
>From 31414b813fa0d9fe50a2e1b0dfcb546d1aac3342 Mon Sep 17 00:00:00 2001
From: Victor Mustya <victor.mustya at intel.com>
Date: Mon, 20 Apr 2026 10:34:13 -0700
Subject: [PATCH 2/7] Fix formatting
---
llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp | 3 +--
1 file changed, 1 insertion(+), 2 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index 7414880ecb05b..0172f91e5f108 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -6287,8 +6287,7 @@ bool SPIRVInstructionSelector::selectAbort(MachineInstr &I) const {
case SPIRV::OpTypeMatrix:
case SPIRV::OpTypeArray:
// Operand 1 holds the element/component type id.
- Worklist.push_back(
- GR.getSPIRVTypeForVReg(Ty->getOperand(1).getReg()));
+ Worklist.push_back(GR.getSPIRVTypeForVReg(Ty->getOperand(1).getReg()));
break;
case SPIRV::OpTypeStruct:
// Operands 1..N hold the field type ids.
>From d0da52015eee3869caffbe9f3886148f2b25e9b3 Mon Sep 17 00:00:00 2001
From: Victor Mustya <victor.mustya at intel.com>
Date: Wed, 22 Apr 2026 08:46:18 -0700
Subject: [PATCH 3/7] Support pointers as concrete types; add a debug check for
the abort followed by unreachable
---
llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp | 35 ++++++++++++-------
.../Target/SPIRV/SPIRVInstructionSelector.cpp | 14 ++++----
.../abort-invalid-message-type.ll | 35 -------------------
.../extensions/SPV_KHR_abort/abort-pointer.ll | 24 +++++++++++++
4 files changed, 53 insertions(+), 55 deletions(-)
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-pointer.ll
diff --git a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
index 6e109c1dd74ab..f0b95d418f249 100644
--- a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
@@ -2253,20 +2253,29 @@ Instruction *SPIRVEmitIntrinsics::visitAtomicCmpXchgInst(AtomicCmpXchgInst &I) {
Instruction *SPIRVEmitIntrinsics::visitUnreachableInst(UnreachableInst &I) {
IRBuilder<> B(I.getParent());
B.SetInsertPoint(&I);
- // OpAbortKHR is itself a SPIR-V block terminator. If the previous instruction
- // is a call to llvm.spv.abort, do not emit an additional OpUnreachable, which
- // would leave the block with two terminators and produce invalid SPIR-V.
- for (Instruction *Prev = I.getPrevNode(); Prev; Prev = Prev->getPrevNode()) {
- if (Prev->isDebugOrPseudoInst())
- continue;
- auto *CI = dyn_cast<CallInst>(Prev);
- if (!CI)
- break;
- Intrinsic::ID IID = CI->getIntrinsicID();
- if (IID == Intrinsic::spv_abort)
+ // OpAbortKHR is itself a SPIR-V block terminator. If the immediately
+ // preceding instruction is a call to llvm.spv.abort, do not emit an
+ // additional OpUnreachable, which would leave the SPIR-V block with two
+ // terminators and produce invalid SPIR-V. The check is intentionally limited
+ // to the directly-preceding non-debug instruction: any real instruction
+ // sitting between `llvm.spv.abort` and `unreachable` would also be invalid
+ // SPIR-V (nothing can follow OpAbortKHR in the same block), so assert that
+ // shape if we see an `spv_abort` anywhere earlier in the block.
+ Instruction *Prev = I.getPrevNode();
+ while (Prev && Prev->isDebugOrPseudoInst())
+ Prev = Prev->getPrevNode();
+ if (auto *CI = dyn_cast_or_null<CallInst>(Prev);
+ CI && CI->getIntrinsicID() == Intrinsic::spv_abort)
return &I;
- break;
- }
+#ifndef NDEBUG
+ for (Instruction *P = I.getPrevNode(); P; P = P->getPrevNode()) {
+ auto *CI = dyn_cast<CallInst>(P);
+ if (CI && CI->getIntrinsicID() == Intrinsic::spv_abort)
+ llvm_unreachable("llvm.spv.abort must be the last non-debug instruction "
+ "before its block's `unreachable`; OpAbortKHR is itself "
+ "a SPIR-V block terminator");
+ }
+#endif
B.CreateIntrinsic(Intrinsic::spv_unreachable, {});
return &I;
}
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index 0172f91e5f108..cf15f98aa2dcd 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -6270,18 +6270,18 @@ bool SPIRVInstructionSelector::selectAbort(MachineInstr &I) const {
assert(MsgType && "Message argument of llvm.spv.abort has no SPIR-V type");
// SPV_KHR_abort requires Message Type to be a concrete type. Per the
// SPIR-V "Concrete Type" definition, that means a numerical scalar
- // (int/float), a vector, matrix, or any aggregate (array/struct)
- // recursively containing only such types. OpTypeBool, OpTypeVoid,
- // pointers, opaque handles and similar abstract/non-concrete types are
+ // (int/float), a (physical) pointer, a vector, matrix, or any aggregate
+ // (array/struct) recursively containing only such types. OpTypeBool,
+ // OpTypeVoid, opaque handles and similar abstract/non-concrete types are
// rejected up front rather than emitting invalid SPIR-V. Validate
- // recursively so that e.g. a struct containing a bool or pointer is also
- // rejected.
+ // recursively so that e.g. a struct containing a bool is also rejected.
SmallVector<SPIRVTypeInst, 4> Worklist{MsgType};
while (!Worklist.empty()) {
SPIRVTypeInst Ty = Worklist.pop_back_val();
switch (Ty->getOpcode()) {
case SPIRV::OpTypeInt:
case SPIRV::OpTypeFloat:
+ case SPIRV::OpTypePointer:
break;
case SPIRV::OpTypeVector:
case SPIRV::OpTypeMatrix:
@@ -6297,8 +6297,8 @@ bool SPIRVInstructionSelector::selectAbort(MachineInstr &I) const {
break;
default:
report_fatal_error("llvm.spv.abort message type must be a concrete "
- "SPIR-V type (numerical scalar, vector, matrix, "
- "or aggregate of such types)",
+ "SPIR-V type (numerical scalar, pointer, vector, "
+ "matrix, or aggregate of such types)",
false);
}
}
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-invalid-message-type.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-invalid-message-type.ll
index d491ac70d80de..32eaa905eee53 100644
--- a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-invalid-message-type.ll
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-invalid-message-type.ll
@@ -1,44 +1,9 @@
; RUN: split-file %s %t
-; RUN: not llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %t/top-level.ll -o /dev/null 2>&1 | FileCheck %s
-; RUN: not llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %t/nested-struct.ll -o /dev/null 2>&1 | FileCheck %s
-; RUN: not llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %t/nested-array.ll -o /dev/null 2>&1 | FileCheck %s
; RUN: not llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %t/top-level-bool.ll -o /dev/null 2>&1 | FileCheck %s
; RUN: not llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %t/nested-bool.ll -o /dev/null 2>&1 | FileCheck %s
; CHECK: llvm.spv.abort message type must be a concrete SPIR-V type
-;--- top-level.ll
-declare void @llvm.spv.abort(...) #0
-define void @abort_with_pointer(ptr %p) {
-entry:
- call void (...) @llvm.spv.abort(ptr %p)
- unreachable
-}
-attributes #0 = { noreturn }
-
-;--- nested-struct.ll
-%S = type { i32, ptr }
-declare void @llvm.spv.abort(...) #0
-define void @abort_with_struct_of_pointer(ptr %p) {
-entry:
- %s0 = insertvalue %S poison, i32 0, 0
- %s1 = insertvalue %S %s0, ptr %p, 1
- call void (...) @llvm.spv.abort(%S %s1)
- unreachable
-}
-attributes #0 = { noreturn }
-
-;--- nested-array.ll
-declare void @llvm.spv.abort(...) #0
-define void @abort_with_array_of_pointer(ptr %p) {
-entry:
- %a0 = insertvalue [2 x ptr] poison, ptr %p, 0
- %a1 = insertvalue [2 x ptr] %a0, ptr %p, 1
- call void (...) @llvm.spv.abort([2 x ptr] %a1)
- unreachable
-}
-attributes #0 = { noreturn }
-
;--- top-level-bool.ll
declare void @llvm.spv.abort(...) #0
define void @abort_with_bool(i1 %b) {
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-pointer.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-pointer.ll
new file mode 100644
index 0000000000000..f4790446950b2
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-pointer.ll
@@ -0,0 +1,24 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - -filetype=obj | spirv-val %}
+
+;; Pointers are concrete SPIR-V types and are valid Message Type operands
+;; for OpAbortKHR.
+
+; CHECK-DAG: OpCapability AbortKHR
+; CHECK-DAG: OpExtension "SPV_KHR_abort"
+; CHECK-DAG: %[[#I8:]] = OpTypeInt 8 0
+; CHECK-DAG: %[[#PTR:]] = OpTypePointer CrossWorkgroup %[[#I8]]
+
+; CHECK: OpAbortKHR %[[#PTR]] %{{[0-9]+}}
+; CHECK-NOT: OpUnreachable
+
+declare void @llvm.spv.abort(...) #0
+
+define spir_kernel void @abort_with_pointer(ptr addrspace(1) %p) {
+entry:
+ call void (...) @llvm.spv.abort(ptr addrspace(1) %p)
+ unreachable
+}
+
+attributes #0 = { noreturn }
>From 4936a114561c5458370e4318409854494a8d3da6 Mon Sep 17 00:00:00 2001
From: Victor Mustya <victor.mustya at intel.com>
Date: Wed, 22 Apr 2026 08:48:11 -0700
Subject: [PATCH 4/7] Fix formatting
---
llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
index f0b95d418f249..1de393e74f4f9 100644
--- a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
@@ -2266,7 +2266,7 @@ Instruction *SPIRVEmitIntrinsics::visitUnreachableInst(UnreachableInst &I) {
Prev = Prev->getPrevNode();
if (auto *CI = dyn_cast_or_null<CallInst>(Prev);
CI && CI->getIntrinsicID() == Intrinsic::spv_abort)
- return &I;
+ return &I;
#ifndef NDEBUG
for (Instruction *P = I.getPrevNode(); P; P = P->getPrevNode()) {
auto *CI = dyn_cast<CallInst>(P);
>From 6bacaf143c6e2cb849a0b2bddb84a79a36f3cf49 Mon Sep 17 00:00:00 2001
From: Victor Mustya <victor.mustya at intel.com>
Date: Thu, 30 Apr 2026 07:58:47 -0700
Subject: [PATCH 5/7] Fix issues/suggestions from review; support trap and
ubsantrap intrinsics; add more tests
---
llvm/include/llvm/IR/IntrinsicsSPIRV.td | 2 +-
llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp | 77 +++++++---
.../Target/SPIRV/SPIRVInstructionSelector.cpp | 110 ++++++++------
.../Target/SPIRV/SPIRVPrepareFunctions.cpp | 134 ++++++++++++------
.../abort-composite-construct.ll | 4 +-
.../SPV_KHR_abort/abort-composite.ll | 7 +-
.../SPV_KHR_abort/abort-conditional.ll | 37 +++++
.../SPV_KHR_abort/abort-in-kernel.ll | 58 ++++++++
.../abort-invalid-message-type.ll | 8 +-
.../SPV_KHR_abort/abort-multiple-blocks.ll | 47 ++++++
.../extensions/SPV_KHR_abort/abort-pointer.ll | 4 +-
.../abort-post-terminator-suppression.ll | 68 +++++++++
.../SPIRV/extensions/SPV_KHR_abort/abort.ll | 4 +-
.../SPV_KHR_abort/debugtrap-not-translated.ll | 24 ++++
.../SPV_KHR_abort/no-abort-unaffected.ll | 59 ++++++++
.../extensions/SPV_KHR_abort/trap-basic.ll | 30 ++++
.../SPV_KHR_abort/trap-conditional.ll | 35 +++++
.../SPV_KHR_abort/trap-ext-disabled.ll | 18 +++
.../SPV_KHR_abort/trap-in-kernel.ll | 54 +++++++
.../SPV_KHR_abort/trap-multiple-blocks.ll | 43 ++++++
.../trap-post-terminator-suppression.ll | 55 +++++++
.../SPV_KHR_abort/ubsantrap-basic.ll | 25 ++++
22 files changed, 781 insertions(+), 122 deletions(-)
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-conditional.ll
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-in-kernel.ll
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-multiple-blocks.ll
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-post-terminator-suppression.ll
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/debugtrap-not-translated.ll
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/no-abort-unaffected.ll
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-basic.ll
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-conditional.ll
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-ext-disabled.ll
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-in-kernel.ll
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-multiple-blocks.ll
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-post-terminator-suppression.ll
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/ubsantrap-basic.ll
diff --git a/llvm/include/llvm/IR/IntrinsicsSPIRV.td b/llvm/include/llvm/IR/IntrinsicsSPIRV.td
index e3d5eda860fcb..5fef26d859d03 100644
--- a/llvm/include/llvm/IR/IntrinsicsSPIRV.td
+++ b/llvm/include/llvm/IR/IntrinsicsSPIRV.td
@@ -41,7 +41,7 @@ let TargetPrefix = "spv" in {
def int_spv_selection_merge : Intrinsic<[], [llvm_any_ty, llvm_i32_ty], [ImmArg<ArgIndex<1>>]>;
def int_spv_cmpxchg : Intrinsic<[llvm_i32_ty], [llvm_any_ty, llvm_vararg_ty]>;
def int_spv_unreachable : Intrinsic<[], []>;
- def int_spv_abort : Intrinsic<[], [llvm_vararg_ty], [IntrNoReturn]>;
+ def int_spv_abort : Intrinsic<[], [llvm_any_ty], [IntrNoReturn]>;
def int_spv_alloca : Intrinsic<[llvm_any_ty], [llvm_i32_ty], [ImmArg<ArgIndex<0>>]>;
def int_spv_alloca_array : Intrinsic<[llvm_any_ty], [llvm_anyint_ty, llvm_i32_ty], [ImmArg<ArgIndex<1>>]>;
def int_spv_undef : Intrinsic<[llvm_i32_ty], []>;
diff --git a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
index 8d91ba06b5aee..39432cf193a93 100644
--- a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
@@ -1542,6 +1542,22 @@ void SPIRVEmitIntrinsics::replaceMemInstrUses(Instruction *Old,
} else if (isMemInstrToReplace(U) || isa<ReturnInst>(U) ||
isa<CallInst>(U)) {
U->replaceUsesOfWith(Old, New);
+ // For a `llvm.spv.abort` call whose composite message argument was
+ // rewritten to a value-id (i32), also retarget the call to a matching
+ // intrinsic declaration so the IR verifier is satisfied. The SPIR-V
+ // type of the value is tracked via the GlobalRegistry, so the selector
+ // still emits OpAbortKHR with the original composite type.
+ if (auto *CI = dyn_cast<CallInst>(U);
+ CI && CI->getIntrinsicID() == Intrinsic::spv_abort) {
+ Type *NewArgTy = New->getType();
+ Type *ExpectedArgTy = CI->getFunctionType()->getParamType(0);
+ if (NewArgTy != ExpectedArgTy) {
+ Module *M = CI->getModule();
+ Function *NewF = Intrinsic::getOrInsertDeclaration(
+ M, Intrinsic::spv_abort, {NewArgTy});
+ CI->setCalledFunction(NewF);
+ }
+ }
} else if (auto *Phi = dyn_cast<PHINode>(U)) {
if (Phi->getType() != New->getType()) {
Phi->mutateType(New->getType());
@@ -2343,32 +2359,47 @@ Instruction *SPIRVEmitIntrinsics::visitAtomicCmpXchgInst(AtomicCmpXchgInst &I) {
return NewI;
}
-Instruction *SPIRVEmitIntrinsics::visitUnreachableInst(UnreachableInst &I) {
- IRBuilder<> B(I.getParent());
- B.SetInsertPoint(&I);
- // OpAbortKHR is itself a SPIR-V block terminator. If the immediately
- // preceding instruction is a call to llvm.spv.abort, do not emit an
- // additional OpUnreachable, which would leave the SPIR-V block with two
- // terminators and produce invalid SPIR-V. The check is intentionally limited
- // to the directly-preceding non-debug instruction: any real instruction
- // sitting between `llvm.spv.abort` and `unreachable` would also be invalid
- // SPIR-V (nothing can follow OpAbortKHR in the same block), so assert that
- // shape if we see an `spv_abort` anywhere earlier in the block.
- Instruction *Prev = I.getPrevNode();
+static bool isAbortCall(const Instruction &I, const SPIRVSubtarget &ST) {
+ auto *CI = dyn_cast<CallInst>(&I);
+ if (!CI)
+ return false;
+ switch (CI->getIntrinsicID()) {
+ case Intrinsic::spv_abort:
+ return true;
+ case Intrinsic::trap:
+ case Intrinsic::ubsantrap:
+ // When the extension is enabled, selection lowers these to OpAbortKHR.
+ return ST.canUseExtension(SPIRV::Extension::SPV_KHR_abort);
+ default:
+ return false;
+ }
+}
+
+// The OpAbortKHR instruction itself is a block terminator, so we don't need to
+// emit an extra OpUnreachable instruction.
+static bool precededByAbortIntrinsic(const UnreachableInst &I,
+ const SPIRVSubtarget &ST) {
+ // Find a previous non-debug instruction.
+ const Instruction *Prev = I.getPrevNode();
while (Prev && Prev->isDebugOrPseudoInst())
Prev = Prev->getPrevNode();
- if (auto *CI = dyn_cast_or_null<CallInst>(Prev);
- CI && CI->getIntrinsicID() == Intrinsic::spv_abort)
+
+ if (Prev && isAbortCall(*Prev, ST))
+ return true;
+
+ assert(llvm::none_of(
+ *I.getParent(),
+ [&ST](const Instruction &II) { return isAbortCall(II, ST); }) &&
+ "abort-like call must be the last non-debug instruction before its "
+ "block's terminator");
+ return false;
+}
+
+Instruction *SPIRVEmitIntrinsics::visitUnreachableInst(UnreachableInst &I) {
+ const SPIRVSubtarget &ST = TM.getSubtarget<SPIRVSubtarget>(*I.getFunction());
+ if (precededByAbortIntrinsic(I, ST))
return &I;
-#ifndef NDEBUG
- for (Instruction *P = I.getPrevNode(); P; P = P->getPrevNode()) {
- auto *CI = dyn_cast<CallInst>(P);
- if (CI && CI->getIntrinsicID() == Intrinsic::spv_abort)
- llvm_unreachable("llvm.spv.abort must be the last non-debug instruction "
- "before its block's `unreachable`; OpAbortKHR is itself "
- "a SPIR-V block terminator");
- }
-#endif
+ IRBuilder<> B(&I);
B.CreateIntrinsic(Intrinsic::spv_unreachable, {});
return &I;
}
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index 48728f44f2e2b..0e7d790d4340f 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -349,6 +349,7 @@ class SPIRVInstructionSelector : public InstructionSelector {
bool diagnoseUnsupported(const MachineInstr &I, const Twine &Msg) const;
bool selectAbort(MachineInstr &I) const;
+ bool selectTrap(MachineInstr &I) const;
bool selectFrameIndex(Register ResVReg, SPIRVTypeInst ResType,
MachineInstr &I) const;
bool selectAllocaArray(Register ResVReg, SPIRVTypeInst ResType,
@@ -1376,11 +1377,13 @@ bool SPIRVInstructionSelector::spvSelect(Register ResVReg,
case TargetOpcode::G_UNMERGE_VALUES:
return selectUnmergeValues(I);
+ case TargetOpcode::G_TRAP:
+ case TargetOpcode::G_UBSANTRAP:
+ return selectTrap(I);
+
// Discard gen opcodes for intrinsics which we do not expect to actually
// represent code after lowering or intrinsics which are not implemented but
// should not crash when found in a customer's LLVM IR input.
- case TargetOpcode::G_TRAP:
- case TargetOpcode::G_UBSANTRAP:
case TargetOpcode::DBG_LABEL:
return true;
case TargetOpcode::G_DEBUGTRAP:
@@ -1740,9 +1743,9 @@ bool SPIRVInstructionSelector::selectPopCount(Register ResVReg,
SPIRVTypeInst ResType,
MachineInstr &I,
unsigned Opcode) const {
- // Vulkan restricts OpBitCount to 32-bit integers or vectors of 32-bit
- // integers unless VK_KHR_maintenance9 is enabled. Until VK_KHR_maintaince9
- // is core we will not generate OpBitCount with any other types when
+ // Vulkan restricts OpBitCount to 32-bit integers or vectors of 32-bit
+ // integers unless VK_KHR_maintenance9 is enabled. Until VK_KHR_maintaince9
+ // is core we will not generate OpBitCount with any other types when
// targeting Vulkan.
if (!STI.getTargetTriple().isVulkanOS())
return selectUnOp(ResVReg, ResType, I, Opcode);
@@ -6329,55 +6332,78 @@ bool SPIRVInstructionSelector::selectAllocaArray(Register ResVReg,
return true;
}
-bool SPIRVInstructionSelector::selectAbort(MachineInstr &I) const {
- if (!STI.canUseExtension(SPIRV::Extension::SPV_KHR_abort))
- report_fatal_error("OpAbortKHR instruction requires the following "
- "SPIR-V extension: SPV_KHR_abort",
- false);
- // The intrinsic is declared as variadic so it can carry composite message
- // types (vectors and structs) without LLVM IR mangling restrictions, but
- // OpAbortKHR takes exactly one Message operand.
- if (I.getNumExplicitOperands() != 2)
- report_fatal_error("llvm.spv.abort must be called with exactly one "
- "message argument",
- false);
- Register MsgReg = I.getOperand(1).getReg();
- SPIRVTypeInst MsgType = GR.getSPIRVTypeForVReg(MsgReg);
- assert(MsgType && "Message argument of llvm.spv.abort has no SPIR-V type");
- // SPV_KHR_abort requires Message Type to be a concrete type. Per the
- // SPIR-V "Concrete Type" definition, that means a numerical scalar
- // (int/float), a (physical) pointer, a vector, matrix, or any aggregate
- // (array/struct) recursively containing only such types. OpTypeBool,
- // OpTypeVoid, opaque handles and similar abstract/non-concrete types are
- // rejected up front rather than emitting invalid SPIR-V. Validate
- // recursively so that e.g. a struct containing a bool is also rejected.
- SmallVector<SPIRVTypeInst, 4> Worklist{MsgType};
+// Returns true iff `Ty` is a concrete SPIR-V type per the SPV_KHR_abort
+// definition: a numerical scalar (int/float), a (physical) pointer, a vector,
+// matrix or any aggregate (array/struct) recursively containing only such
+// types. OpTypeBool, OpTypeVoid, opaque handles and similar abstract
+// non-concrete types are rejected.
+static bool isConcreteSPIRVType(SPIRVTypeInst Ty,
+ const SPIRVGlobalRegistry &GR) {
+ SmallVector<SPIRVTypeInst, 4> Worklist{Ty};
while (!Worklist.empty()) {
- SPIRVTypeInst Ty = Worklist.pop_back_val();
- switch (Ty->getOpcode()) {
+ SPIRVTypeInst T = Worklist.pop_back_val();
+ switch (T->getOpcode()) {
case SPIRV::OpTypeInt:
case SPIRV::OpTypeFloat:
case SPIRV::OpTypePointer:
break;
case SPIRV::OpTypeVector:
case SPIRV::OpTypeMatrix:
- case SPIRV::OpTypeArray:
- // Operand 1 holds the element/component type id.
- Worklist.push_back(GR.getSPIRVTypeForVReg(Ty->getOperand(1).getReg()));
- break;
+ case SPIRV::OpTypeArray: {
+ Register OperandReg = T->getOperand(1).getReg();
+ SPIRVTypeInst ElementT = GR.getSPIRVTypeForVReg(OperandReg);
+ Worklist.push_back(ElementT);
+ } break;
case SPIRV::OpTypeStruct:
- // Operands 1..N hold the field type ids.
- for (unsigned Idx = 1, E = Ty->getNumOperands(); Idx < E; ++Idx)
- Worklist.push_back(
- GR.getSPIRVTypeForVReg(Ty->getOperand(Idx).getReg()));
+ for (unsigned Idx = 1, E = T->getNumOperands(); Idx < E; ++Idx) {
+ Register OperandReg = T->getOperand(Idx).getReg();
+ SPIRVTypeInst ElementT = GR.getSPIRVTypeForVReg(OperandReg);
+ Worklist.push_back(ElementT);
+ }
break;
default:
- report_fatal_error("llvm.spv.abort message type must be a concrete "
- "SPIR-V type (numerical scalar, pointer, vector, "
- "matrix, or aggregate of such types)",
- false);
+ return false;
}
}
+ return true;
+}
+
+bool SPIRVInstructionSelector::selectAbort(MachineInstr &I) const {
+ assert(I.getNumExplicitOperands() == 2);
+
+ Register MsgReg = I.getOperand(1).getReg();
+ SPIRVTypeInst MsgType = GR.getSPIRVTypeForVReg(MsgReg);
+ assert(MsgType && "Message argument of llvm.spv.abort has no SPIR-V type");
+
+ if (!isConcreteSPIRVType(MsgType, GR))
+ return diagnoseUnsupported(
+ I,
+ "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
+ "scalar, pointer, vector, matrix, or aggregate of such types)");
+
+ MachineBasicBlock &BB = *I.getParent();
+ BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpAbortKHR))
+ .addUse(GR.getSPIRVTypeID(MsgType))
+ .addUse(MsgReg)
+ .constrainAllUses(TII, TRI, RBI);
+ return true;
+}
+
+bool SPIRVInstructionSelector::selectTrap(MachineInstr &I) const {
+ // When the SPV_KHR_abort extension is disabled, drop the G_TRAP and
+ // G_UBSANTRAP silently.
+ if (!STI.canUseExtension(SPIRV::Extension::SPV_KHR_abort))
+ return true;
+
+ // Use the 32-bit integer constant for the abort "message" argument:
+ // - G_UBSANTRAP operand is zero-extended to 32 bits.
+ // - "All ones" constant is used for G_TRAP.
+ uint32_t MsgVal = ~0u;
+ if (I.getOpcode() == TargetOpcode::G_UBSANTRAP)
+ MsgVal = static_cast<uint32_t>(I.getOperand(0).getImm());
+
+ SPIRVTypeInst MsgType = GR.getOrCreateSPIRVIntegerType(32, I, TII);
+ Register MsgReg = buildI32Constant(MsgVal, I, MsgType);
MachineBasicBlock &BB = *I.getParent();
BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpAbortKHR))
.addUse(GR.getSPIRVTypeID(MsgType))
diff --git a/llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp b/llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp
index 6dfa1b6aa812f..cf079d2bad409 100644
--- a/llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp
@@ -47,6 +47,7 @@ class SPIRVPrepareFunctionsImpl {
const SPIRVTargetMachine &TM;
bool substituteIntrinsicCalls(Function *F);
bool substituteAbortKHRCalls(Function *F);
+ bool terminateBlocksAfterTrap(Module &M, Intrinsic::ID IID);
Function *removeAggregateTypesFromSignature(Function *F);
bool removeAggregateTypesFromCalls(Function *F);
@@ -543,8 +544,10 @@ addFunctionTypeMutation(NamedMDNode *NMD,
Function *
SPIRVPrepareFunctionsImpl::removeAggregateTypesFromSignature(Function *F) {
bool IsRetAggr = F->getReturnType()->isAggregateType();
- // Allow intrinsics with aggregate return type to reach GlobalISel
- if (F->isIntrinsic() && IsRetAggr)
+ // Allow intrinsics with aggregate return/argument types to reach GlobalISel.
+ // Renaming/mutating the signature of an intrinsic would desync its name from
+ // its argument types and break the IR verifier.
+ if (F->isIntrinsic())
return F;
IRBuilder<> B(F->getContext());
@@ -608,60 +611,102 @@ SPIRVPrepareFunctionsImpl::removeAggregateTypesFromSignature(Function *F) {
return NewF;
}
-// Replace OpenCL/SPIR-V style calls to `__spirv_AbortKHR(message)` with calls
-// to the `llvm.spv.abort` target intrinsic, so that they go through the same
-// instruction-selection path as the intrinsic and get lowered to OpAbortKHR.
+// Returns true iff `F`'s name resolves (after OpenCL/SPIR-V demangling and
+// builtin-name lookup) to the SPIR-V friendly built-in `__spirv_AbortKHR`.
+static bool isAbortKHRBuiltin(const Function &F) {
+ if (F.isIntrinsic())
+ return false;
+ StringRef Name = F.getName();
+ std::string Demangled = getOclOrSpirvBuiltinDemangledName(Name);
+ if (Demangled.empty())
+ return false;
+ return SPIRV::lookupBuiltinNameHelper(Demangled) == "__spirv_AbortKHR";
+}
+
+// Rewrites a single call to `__spirv_AbortKHR` into a call to the
+// `llvm.spv.abort` target intrinsic, then re-terminates the block with
+// `unreachable`. OpAbortKHR is itself a SPIR-V function-termination
+// instruction and must be the last instruction in its block, so any trailing
+// stores/lifetime intrinsics/`ret` emitted by the OpenCL ABI are dropped.
+// `changeToUnreachable` cleans up any successor PHI predecessor entries.
+static void rewriteAbortKHRCall(CallInst *CI) {
+ IRBuilder<> B(CI);
+ Value *Msg = CI->getArgOperand(0);
+ // The OpenCL C ABI may pass aggregate arguments by pointer (byval). In that
+ // case load the underlying value so that OpAbortKHR receives the composite
+ // itself, as required by the SPV_KHR_abort spec ("Message Type must be a
+ // concrete type").
+ if (CI->isByValArgument(0)) {
+ Type *AggTy = CI->getParamByValType(0);
+ Msg = B.CreateLoad(AggTy, Msg);
+ }
+ B.CreateIntrinsic(Intrinsic::spv_abort, {Msg->getType()}, {Msg});
+ changeToUnreachable(CI);
+}
+
+// Replace OpenCL/SPIR-V style calls to `__spirv_AbortKHR(message)` (i.e.
+// calls to `F` when `F` is the `__spirv_AbortKHR` built-in) with calls to the
+// `llvm.spv.abort` target intrinsic.
bool SPIRVPrepareFunctionsImpl::substituteAbortKHRCalls(Function *F) {
- if (F->isDeclaration())
+ if (!isAbortKHRBuiltin(*F))
return false;
SmallVector<CallInst *> Calls;
- for (Instruction &I : instructions(F)) {
- auto *CI = dyn_cast<CallInst>(&I);
- if (!CI)
- continue;
- Function *Callee = CI->getCalledFunction();
- if (!Callee || Callee->isIntrinsic())
- continue;
- StringRef Demangled = Callee->getName();
- std::string DemangledStr = getOclOrSpirvBuiltinDemangledName(Demangled);
- if (DemangledStr.empty())
- continue;
- std::string BuiltinName = SPIRV::lookupBuiltinNameHelper(DemangledStr);
- if (StringRef(BuiltinName) != "__spirv_AbortKHR")
+ for (User *U : F->users()) {
+ auto *CI = dyn_cast<CallInst>(U);
+ if (!CI || CI->getCalledFunction() != F)
continue;
if (CI->arg_size() != 1)
continue;
Calls.push_back(CI);
}
- if (Calls.empty())
+ for (CallInst *CI : Calls)
+ rewriteAbortKHRCall(CI);
+
+ return !Calls.empty();
+}
+
+// When the SPV_KHR_abort extension is enabled, `llvm.trap` and
+// `llvm.ubsantrap` are lowered to `OpAbortKHR` during instruction selection.
+// `OpAbortKHR` is itself a SPIR-V block terminator, so any instructions that
+// follow the trap call within the same basic block (e.g. `ret`, lifetime
+// markers) would produce SPIR-V ops after `OpAbortKHR` and break validation.
+// Terminate the block right after each call to the trap intrinsics by replacing
+// the next instruction with `unreachable`.
+bool SPIRVPrepareFunctionsImpl::terminateBlocksAfterTrap(Module &M,
+ Intrinsic::ID IID) {
+ assert((IID == Intrinsic::trap || IID == Intrinsic::ubsantrap) &&
+ "Expected trap intrinsic ID");
+
+ Function *F = Intrinsic::getDeclarationIfExists(&M, IID);
+ if (!F)
+ return false;
+
+ // If the target doesn't support SPV_KHR_abort, we won't be able to lower
+ // the trap intrinsic to OpAbortKHR, so we can skip the block-terminating
+ // transformation.
+ const auto &ST = TM.getSubtarget<SPIRVSubtarget>(*F);
+ if (!ST.canUseExtension(SPIRV::Extension::SPV_KHR_abort))
return false;
+ SmallVector<CallInst *> Calls;
+ for (User *U : F->users()) {
+ auto *CI = dyn_cast<CallInst>(U);
+ if (!CI || CI->getCalledFunction() != F)
+ continue;
+ Calls.push_back(CI);
+ }
+
+ bool Changed = false;
for (CallInst *CI : Calls) {
- IRBuilder<> B(CI);
- Value *Msg = CI->getArgOperand(0);
- // The OpenCL C ABI may pass aggregate arguments by pointer (byval). In
- // that case load the underlying value so that OpAbortKHR receives the
- // composite itself, as required by the SPV_KHR_abort spec ("Message Type
- // must be a concrete type").
- if (CI->isByValArgument(0)) {
- Type *AggTy = CI->getParamByValType(0);
- Msg = B.CreateLoad(AggTy, Msg);
- }
- B.CreateIntrinsic(Intrinsic::spv_abort, {}, {Msg});
- // OpAbortKHR is itself a SPIR-V function-termination instruction and must
- // be the last instruction in its block. Drop the original call and
- // everything that follows it (the OpenCL ABI typically appends stores into
- // the return slot and a `ret`), and re-terminate the block with
- // `unreachable`. We use changeToUnreachable so that any successor PHI
- // nodes have the now-removed predecessor edge cleaned up; otherwise the
- // IR verifier would reject mismatched PHI incoming entries. The matching
- // suppression in SPIRVEmitIntrinsics::visitUnreachableInst ensures no
- // extra OpUnreachable is emitted after OpAbortKHR.
- changeToUnreachable(CI);
+ Instruction *Next = CI->getNextNode();
+ if (!Next || isa<UnreachableInst>(Next))
+ continue;
+ changeToUnreachable(Next);
+ Changed = true;
}
- return true;
+ return Changed;
}
static std::string fixMultiOutputConstraintString(StringRef Constraints) {
@@ -778,9 +823,12 @@ bool SPIRVPrepareFunctionsImpl::runOnModule(Module &M) {
Changed = true;
}
+ Changed |= terminateBlocksAfterTrap(M, Intrinsic::trap);
+ Changed |= terminateBlocksAfterTrap(M, Intrinsic::ubsantrap);
+
for (Function &F : M) {
- Changed |= substituteIntrinsicCalls(&F);
Changed |= substituteAbortKHRCalls(&F);
+ Changed |= substituteIntrinsicCalls(&F);
Changed |= sortBlocks(F);
Changed |= removeAggregateTypesFromCalls(&F);
}
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-composite-construct.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-composite-construct.ll
index b7ab99b47b0cd..9bea2744ac281 100644
--- a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-composite-construct.ll
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-composite-construct.ll
@@ -15,7 +15,7 @@
; CHECK: OpAbortKHR %[[#V2]] %[[#CC]]
; CHECK-NOT: OpUnreachable
-declare void @llvm.spv.abort(...) #0
+declare void @llvm.spv.abort(<2 x i32>) #0
define spir_kernel void @abort_composite_construct(i32 %a, i32 %b) {
entry:
@@ -23,7 +23,7 @@ entry:
%vb = insertelement <1 x i32> poison, i32 %b, i32 0
%v = shufflevector <1 x i32> %va, <1 x i32> %vb,
<2 x i32> <i32 0, i32 1>
- call void (...) @llvm.spv.abort(<2 x i32> %v)
+ call void @llvm.spv.abort(<2 x i32> %v)
unreachable
}
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-composite.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-composite.ll
index f8c779dc7b891..fd352bffd87cd 100644
--- a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-composite.ll
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-composite.ll
@@ -27,14 +27,15 @@
%struct.Msg = type { i32, i32, float }
-declare void @llvm.spv.abort(...) #0
+declare void @_Z16__spirv_AbortKHR3Msg(%struct.Msg) #0
+declare void @_Z16__spirv_AbortKHRDv4_j(<4 x i32>) #0
define spir_kernel void @abort_with_struct(i32 %x, i32 %y, float %z) {
entry:
%m0 = insertvalue %struct.Msg poison, i32 %x, 0
%m1 = insertvalue %struct.Msg %m0, i32 %y, 1
%m2 = insertvalue %struct.Msg %m1, float %z, 2
- call void (...) @llvm.spv.abort(%struct.Msg %m2)
+ call void @_Z16__spirv_AbortKHR3Msg(%struct.Msg %m2)
unreachable
}
@@ -44,7 +45,7 @@ entry:
%v1 = insertelement <4 x i32> %v0, i32 %b, i32 1
%v2 = insertelement <4 x i32> %v1, i32 %c, i32 2
%v3 = insertelement <4 x i32> %v2, i32 %d, i32 3
- call void (...) @llvm.spv.abort(<4 x i32> %v3)
+ call void @_Z16__spirv_AbortKHRDv4_j(<4 x i32> %v3)
unreachable
}
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-conditional.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-conditional.ll
new file mode 100644
index 0000000000000..60bf9648f93d0
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-conditional.ll
@@ -0,0 +1,37 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - -filetype=obj | spirv-val %}
+
+;; Conditional abort -- the assert() pattern where only some paths abort.
+;; Verifies that:
+;; 1. The abort BB ends with OpAbortKHR (no OpUnreachable after it).
+;; 2. The non-abort BB is unaffected (still has OpReturn).
+;; 3. No double terminators in the abort block.
+
+; CHECK-DAG: OpCapability AbortKHR
+; CHECK-DAG: OpExtension "SPV_KHR_abort"
+
+; CHECK-DAG: OpBranchConditional
+; CHECK-DAG: OpReturn
+; CHECK-DAG: OpAbortKHR
+
+;; Crucially, no OpUnreachable should appear: the trap block's `unreachable`
+;; is consumed by OpAbortKHR.
+; CHECK-NOT: OpUnreachable
+
+declare spir_func void @_Z16__spirv_AbortKHRj(i32) #0
+
+define spir_func void @assert_like(i32 %gid, i32 %N, i32 %msg) {
+entry:
+ %cmp = icmp slt i32 %gid, %N
+ br i1 %cmp, label %ok, label %trap
+
+ok:
+ ret void
+
+trap:
+ call spir_func void @_Z16__spirv_AbortKHRj(i32 %msg)
+ unreachable
+}
+
+attributes #0 = { noreturn }
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-in-kernel.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-in-kernel.ll
new file mode 100644
index 0000000000000..5f9921413657d
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-in-kernel.ll
@@ -0,0 +1,58 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - -filetype=obj | spirv-val %}
+
+;; Kernel entry point invokes a helper that aborts. Models the device library
+;; __assert_fail pattern: the kernel calls __assert_fail_internal which calls
+;; __spirv_AbortKHR. The kernel's call site is followed by `unreachable`, which
+;; must be preserved (the abort is in the callee, not the caller).
+
+; CHECK-DAG: OpCapability AbortKHR
+; CHECK-DAG: OpExtension "SPV_KHR_abort"
+; CHECK-DAG: OpEntryPoint Kernel %{{[0-9]+}} "test_kernel"
+
+;; Helper function: abort lowered to OpAbortKHR.
+; CHECK: OpFunction
+; CHECK: OpAbortKHR
+; CHECK-NEXT: OpFunctionEnd
+
+;; Kernel: conditional branch, then a function call followed by OpUnreachable
+;; in the assert.fail block (the unreachable after the call to the abort
+;; helper is preserved because the helper, not the kernel, contains the
+;; OpAbortKHR).
+; CHECK: OpFunction
+; CHECK-DAG: OpBranchConditional
+; CHECK-DAG: OpReturn
+; CHECK-DAG: OpFunctionCall
+; CHECK-DAG: OpUnreachable
+; CHECK: OpFunctionEnd
+
+declare spir_func void @_Z16__spirv_AbortKHRj(i32) #0
+declare spir_func i64 @_Z13get_global_idj(i32) #1
+
+; Models __assert_fail from device libraries.
+define spir_func void @__assert_fail_internal(i32 %msg) #2 {
+entry:
+ call spir_func void @_Z16__spirv_AbortKHRj(i32 %msg)
+ unreachable
+}
+
+; Kernel entry point with conditional assert.
+define spir_kernel void @test_kernel(ptr addrspace(1) %in, i32 %N) {
+entry:
+ %gid = call spir_func i64 @_Z13get_global_idj(i32 0)
+ %gid32 = trunc i64 %gid to i32
+ %cmp = icmp slt i32 %gid32, %N
+ br i1 %cmp, label %ok, label %assert.fail
+
+ok:
+ ret void
+
+assert.fail:
+ call spir_func void @__assert_fail_internal(i32 42)
+ unreachable
+}
+
+attributes #0 = { noreturn }
+attributes #1 = { nounwind }
+attributes #2 = { noinline noreturn nounwind }
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-invalid-message-type.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-invalid-message-type.ll
index 32eaa905eee53..f5c2dcc34d0c5 100644
--- a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-invalid-message-type.ll
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-invalid-message-type.ll
@@ -5,22 +5,22 @@
; CHECK: llvm.spv.abort message type must be a concrete SPIR-V type
;--- top-level-bool.ll
-declare void @llvm.spv.abort(...) #0
+declare void @llvm.spv.abort.i1(i1) #0
define void @abort_with_bool(i1 %b) {
entry:
- call void (...) @llvm.spv.abort(i1 %b)
+ call void @llvm.spv.abort.i1(i1 %b)
unreachable
}
attributes #0 = { noreturn }
;--- nested-bool.ll
%B = type { i32, i1 }
-declare void @llvm.spv.abort(...) #0
+declare void @llvm.spv.abort.s_Bs(%B) #0
define void @abort_with_struct_of_bool(i1 %b) {
entry:
%s0 = insertvalue %B poison, i32 0, 0
%s1 = insertvalue %B %s0, i1 %b, 1
- call void (...) @llvm.spv.abort(%B %s1)
+ call void @llvm.spv.abort.s_Bs(%B %s1)
unreachable
}
attributes #0 = { noreturn }
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-multiple-blocks.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-multiple-blocks.ll
new file mode 100644
index 0000000000000..95d64a06cbf7b
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-multiple-blocks.ll
@@ -0,0 +1,47 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - -filetype=obj | spirv-val %}
+
+;; Multiple basic blocks: some abort, some don't.
+;; Verifies that:
+;; 1. Non-abort blocks are unaffected (normal terminators preserved).
+;; 2. Multiple abort blocks in the same function each get their own OpAbortKHR.
+;; 3. No cross-contamination between blocks.
+
+; CHECK-DAG: OpCapability AbortKHR
+; CHECK-DAG: OpExtension "SPV_KHR_abort"
+
+; CHECK: OpFunction
+; CHECK-DAG: OpBranchConditional
+; CHECK-DAG: OpBranchConditional
+; CHECK-DAG: OpReturnValue
+; CHECK-DAG: OpAbortKHR
+; CHECK-DAG: OpAbortKHR
+; CHECK: OpFunctionEnd
+; CHECK-NOT: OpUnreachable
+
+declare spir_func void @_Z16__spirv_AbortKHRj(i32) #0
+
+define spir_func i32 @multi_abort(i32 %x, i32 %m1, i32 %m2) {
+entry:
+ %cmp1 = icmp sgt i32 %x, 0
+ br i1 %cmp1, label %work, label %err1
+
+work:
+ %result = mul i32 %x, 42
+ %cmp2 = icmp slt i32 %result, 1000
+ br i1 %cmp2, label %ret, label %err2
+
+ret:
+ ret i32 %result
+
+err1:
+ call spir_func void @_Z16__spirv_AbortKHRj(i32 %m1)
+ unreachable
+
+err2:
+ call spir_func void @_Z16__spirv_AbortKHRj(i32 %m2)
+ unreachable
+}
+
+attributes #0 = { noreturn }
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-pointer.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-pointer.ll
index f4790446950b2..3693091759610 100644
--- a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-pointer.ll
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-pointer.ll
@@ -13,11 +13,11 @@
; CHECK: OpAbortKHR %[[#PTR]] %{{[0-9]+}}
; CHECK-NOT: OpUnreachable
-declare void @llvm.spv.abort(...) #0
+declare void @llvm.spv.abort.p1(ptr addrspace(1)) #0
define spir_kernel void @abort_with_pointer(ptr addrspace(1) %p) {
entry:
- call void (...) @llvm.spv.abort(ptr addrspace(1) %p)
+ call void @llvm.spv.abort.p1(ptr addrspace(1) %p)
unreachable
}
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-post-terminator-suppression.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-post-terminator-suppression.ll
new file mode 100644
index 0000000000000..45d4b6cd37101
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-post-terminator-suppression.ll
@@ -0,0 +1,68 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - -filetype=obj | spirv-val %}
+
+;; Edge case: instructions after the __spirv_AbortKHR call in the same basic
+;; block.
+;;
+;; In real code (e.g. device libraries' __assert_fail), the pattern is:
+;; call void @__spirv_AbortKHR(i32 %msg)
+;; ; ... possibly lifetime.end intrinsics ...
+;; ret void ; or unreachable
+;;
+;; OpAbortKHR is itself a SPIR-V block terminator, so all subsequent
+;; instructions in the same BB must be suppressed to produce valid SPIR-V.
+;; This test verifies that no instructions appear after OpAbortKHR in any
+;; function.
+
+; CHECK-DAG: OpCapability AbortKHR
+; CHECK-DAG: OpExtension "SPV_KHR_abort"
+
+;; abort then unreachable.
+; CHECK: OpFunction
+; CHECK: OpAbortKHR
+; CHECK-NEXT: OpFunctionEnd
+
+;; abort then ret void.
+; CHECK: OpFunction
+; CHECK: OpAbortKHR
+; CHECK-NEXT: OpFunctionEnd
+
+;; abort then lifetime.end + ret void.
+; CHECK: OpFunction
+; CHECK: OpAbortKHR
+; CHECK-NEXT: OpFunctionEnd
+
+; CHECK-NOT: OpReturn{{[[:space:]]+}}OpFunctionEnd
+; CHECK-NOT: OpUnreachable
+
+declare spir_func void @_Z16__spirv_AbortKHRj(i32) #0
+declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1
+declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1
+
+; Pattern 1: abort + unreachable.
+define spir_func void @abort_then_unreachable(i32 %msg) {
+entry:
+ call spir_func void @_Z16__spirv_AbortKHRj(i32 %msg)
+ unreachable
+}
+
+; Pattern 2: abort + ret void.
+define spir_func void @abort_then_ret(i32 %msg) {
+entry:
+ call spir_func void @_Z16__spirv_AbortKHRj(i32 %msg)
+ ret void
+}
+
+; Pattern 3: abort + lifetime.end + ret void.
+define spir_func void @abort_then_lifetime_ret(i32 %msg) {
+entry:
+ %buf = alloca i8, align 1
+ call void @llvm.lifetime.start.p0(i64 1, ptr %buf)
+ call spir_func void @_Z16__spirv_AbortKHRj(i32 %msg)
+ call void @llvm.lifetime.end.p0(i64 1, ptr %buf)
+ ret void
+}
+
+attributes #0 = { noreturn }
+attributes #1 = { argmemonly nounwind willreturn }
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort.ll
index a91cb02a36927..da1e20ea66877 100644
--- a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort.ll
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort.ll
@@ -15,10 +15,10 @@
; CHECK: OpAbortKHR %[[#I32]] %[[#MSG]]
; CHECK-NOT: OpUnreachable
-declare void @llvm.spv.abort(...)
+declare void @llvm.spv.abort(i32)
define void @abort_with_int() {
entry:
- call void (...) @llvm.spv.abort(i32 42)
+ call void @llvm.spv.abort(i32 42)
unreachable
}
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/debugtrap-not-translated.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/debugtrap-not-translated.ll
new file mode 100644
index 0000000000000..8ac35926e09b1
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/debugtrap-not-translated.ll
@@ -0,0 +1,24 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+
+;; Negative test: llvm.debugtrap must NOT lower to OpAbortKHR. Only llvm.trap
+;; and llvm.ubsantrap are translated to OpAbortKHR; debugtrap is dropped (no
+;; codegen) and the original `unreachable` terminator is preserved.
+
+; CHECK-NOT: OpCapability AbortKHR
+; CHECK-NOT: OpAbortKHR
+
+; CHECK: OpFunction
+; CHECK: OpLabel
+; CHECK: OpUnreachable
+; CHECK: OpFunctionEnd
+
+define spir_func void @uses_debugtrap() {
+entry:
+ call void @llvm.debugtrap()
+ unreachable
+}
+
+declare void @llvm.debugtrap() #0
+
+attributes #0 = { nounwind }
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/no-abort-unaffected.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/no-abort-unaffected.ll
new file mode 100644
index 0000000000000..17e1a0223c7a8
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/no-abort-unaffected.ll
@@ -0,0 +1,59 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - -filetype=obj | spirv-val %}
+
+;; Sanity check: enabling SPV_KHR_abort does not affect functions that don't
+;; abort. Normal terminators (OpReturn, OpReturnValue, OpBranch,
+;; OpUnreachable) must be preserved.
+
+;; No abort instructions emitted anywhere.
+; CHECK-NOT: OpAbortKHR
+
+;; Normal void return preserved.
+; CHECK: OpFunction
+; CHECK: OpReturn
+; CHECK: OpFunctionEnd
+
+;; Normal value return preserved.
+; CHECK: OpFunction
+; CHECK: OpReturnValue
+; CHECK: OpFunctionEnd
+
+;; Plain unreachable preserved (no abort precedes it).
+; CHECK: OpFunction
+; CHECK: OpUnreachable
+; CHECK: OpFunctionEnd
+
+;; Branch preserved.
+; CHECK: OpFunction
+; CHECK: OpBranchConditional
+; CHECK: OpReturnValue
+; CHECK: OpReturnValue
+; CHECK: OpFunctionEnd
+
+define spir_func void @void_return() {
+entry:
+ ret void
+}
+
+define spir_func i32 @value_return(i32 %x) {
+entry:
+ %r = add i32 %x, 1
+ ret i32 %r
+}
+
+define spir_func void @plain_unreachable() {
+entry:
+ unreachable
+}
+
+define spir_func i32 @branched(i1 %cond, i32 %a, i32 %b) {
+entry:
+ br i1 %cond, label %t, label %f
+
+t:
+ ret i32 %a
+
+f:
+ ret i32 %b
+}
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-basic.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-basic.ll
new file mode 100644
index 0000000000000..f28357bb6b95d
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-basic.ll
@@ -0,0 +1,30 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - -filetype=obj | spirv-val %}
+
+;; Without the extension, llvm.trap is dropped silently (existing behavior).
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-NO-EXT
+
+;; llvm.trap lowers to OpAbortKHR with an all-ones (-1) i32 Message constant.
+
+; CHECK-DAG: OpCapability AbortKHR
+; CHECK-DAG: OpExtension "SPV_KHR_abort"
+; CHECK-DAG: %[[#I32:]] = OpTypeInt 32 0
+; CHECK-DAG: %[[#MSG:]] = OpConstant %[[#I32]] 4294967295
+
+; CHECK: OpAbortKHR %[[#I32]] %[[#MSG]]
+; CHECK-NOT: OpUnreachable
+; CHECK-NEXT: OpFunctionEnd
+
+; CHECK-NO-EXT-NOT: OpCapability AbortKHR
+; CHECK-NO-EXT-NOT: OpAbortKHR
+
+define spir_func void @trap_simple() {
+entry:
+ call void @llvm.trap()
+ unreachable
+}
+
+declare void @llvm.trap() #0
+
+attributes #0 = { cold noreturn nounwind }
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-conditional.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-conditional.ll
new file mode 100644
index 0000000000000..317f22df046b6
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-conditional.ll
@@ -0,0 +1,35 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - -filetype=obj | spirv-val %}
+
+;; Conditional trap: assert() pattern where only one path traps.
+;; Verifies that the trap BB ends with OpAbortKHR (no OpUnreachable) and
+;; the non-trap BB still has OpReturn.
+
+; CHECK-DAG: OpCapability AbortKHR
+; CHECK-DAG: OpExtension "SPV_KHR_abort"
+
+; CHECK: OpFunction
+; CHECK-DAG: OpBranchConditional
+; CHECK-DAG: OpReturn
+; CHECK-DAG: OpAbortKHR
+
+;; The trap block's `unreachable` must be consumed by OpAbortKHR.
+; CHECK-NOT: OpUnreachable
+
+define spir_func void @assert_like(i32 %gid, i32 %N) {
+entry:
+ %cmp = icmp slt i32 %gid, %N
+ br i1 %cmp, label %ok, label %trap
+
+ok:
+ ret void
+
+trap:
+ call void @llvm.trap()
+ unreachable
+}
+
+declare void @llvm.trap() #0
+
+attributes #0 = { cold noreturn nounwind }
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-ext-disabled.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-ext-disabled.ll
new file mode 100644
index 0000000000000..7cf78e8553303
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-ext-disabled.ll
@@ -0,0 +1,18 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+;; When SPV_KHR_abort is not enabled, llvm.trap is dropped (existing behavior),
+;; and no AbortKHR capability/extension/instruction is emitted.
+
+; CHECK-NOT: OpCapability AbortKHR
+; CHECK-NOT: OpExtension "SPV_KHR_abort"
+; CHECK-NOT: OpAbortKHR
+
+define spir_func void @trap_with_ext_disabled() {
+entry:
+ call void @llvm.trap()
+ unreachable
+}
+
+declare void @llvm.trap() #0
+
+attributes #0 = { cold noreturn nounwind }
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-in-kernel.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-in-kernel.ll
new file mode 100644
index 0000000000000..c7874dc0bf576
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-in-kernel.ll
@@ -0,0 +1,54 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - -filetype=obj | spirv-val %}
+
+;; Kernel function with assert-like trap pattern. Models the real-world HIP
+;; assert() use case: kernel calls a helper which calls llvm.trap.
+
+; CHECK-DAG: OpCapability AbortKHR
+; CHECK-DAG: OpExtension "SPV_KHR_abort"
+; CHECK-DAG: OpEntryPoint Kernel %{{[0-9]+}} "test_kernel"
+
+;; __assert_fail_internal: trap → OpAbortKHR.
+; CHECK: OpFunction
+; CHECK: OpAbortKHR
+; CHECK-NEXT: OpFunctionEnd
+
+;; test_kernel: conditional branch + return + function call + unreachable.
+; CHECK: OpFunction
+; CHECK-DAG: OpBranchConditional
+; CHECK-DAG: OpReturn
+; CHECK-DAG: OpFunctionCall
+; CHECK-DAG: OpUnreachable
+; CHECK: OpFunctionEnd
+
+declare spir_func i64 @_Z13get_global_idj(i32) #2
+declare void @llvm.trap() #3
+
+; Models __assert_fail from device libraries.
+define spir_func void @__assert_fail_internal() #0 {
+entry:
+ call void @llvm.trap()
+ unreachable
+}
+
+; Kernel entry point with conditional assert.
+define spir_kernel void @test_kernel(ptr addrspace(1) %in, i32 %N) #1 {
+entry:
+ %gid = call spir_func i64 @_Z13get_global_idj(i32 0)
+ %gid32 = trunc i64 %gid to i32
+ %cmp = icmp slt i32 %gid32, %N
+ br i1 %cmp, label %ok, label %assert.fail
+
+ok:
+ ret void
+
+assert.fail:
+ call spir_func void @__assert_fail_internal()
+ unreachable
+}
+
+attributes #0 = { noinline noreturn nounwind }
+attributes #1 = { nounwind }
+attributes #2 = { nounwind }
+attributes #3 = { cold noreturn nounwind }
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-multiple-blocks.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-multiple-blocks.ll
new file mode 100644
index 0000000000000..4db2e616e0716
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-multiple-blocks.ll
@@ -0,0 +1,43 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - -filetype=obj | spirv-val %}
+
+;; Multiple basic blocks: some trap, some don't.
+
+; CHECK-DAG: OpCapability AbortKHR
+; CHECK-DAG: OpExtension "SPV_KHR_abort"
+
+; CHECK: OpFunction
+; CHECK-DAG: OpBranchConditional
+; CHECK-DAG: OpBranchConditional
+; CHECK-DAG: OpReturnValue
+; CHECK-DAG: OpAbortKHR
+; CHECK-DAG: OpAbortKHR
+; CHECK: OpFunctionEnd
+; CHECK-NOT: OpUnreachable
+
+define spir_func i32 @multi_trap(i32 %x) {
+entry:
+ %cmp1 = icmp sgt i32 %x, 0
+ br i1 %cmp1, label %work, label %err1
+
+work:
+ %result = mul i32 %x, 42
+ %cmp2 = icmp slt i32 %result, 1000
+ br i1 %cmp2, label %ret, label %err2
+
+ret:
+ ret i32 %result
+
+err1:
+ call void @llvm.trap()
+ unreachable
+
+err2:
+ call void @llvm.trap()
+ unreachable
+}
+
+declare void @llvm.trap() #0
+
+attributes #0 = { cold noreturn nounwind }
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-post-terminator-suppression.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-post-terminator-suppression.ll
new file mode 100644
index 0000000000000..3c7e94048423f
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/trap-post-terminator-suppression.ll
@@ -0,0 +1,55 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - -filetype=obj | spirv-val %}
+
+;; Edge case: instructions after llvm.trap in the same basic block must be
+;; suppressed because OpAbortKHR is a SPIR-V block terminator.
+
+; CHECK-DAG: OpCapability AbortKHR
+; CHECK-DAG: OpExtension "SPV_KHR_abort"
+
+;; trap then unreachable.
+; CHECK: OpFunction
+; CHECK: OpAbortKHR
+; CHECK-NEXT: OpFunctionEnd
+
+;; trap then ret void.
+; CHECK: OpFunction
+; CHECK: OpAbortKHR
+; CHECK-NEXT: OpFunctionEnd
+
+;; trap then lifetime.end + ret void.
+; CHECK: OpFunction
+; CHECK: OpAbortKHR
+; CHECK-NEXT: OpFunctionEnd
+
+; CHECK-NOT: OpReturn{{[[:space:]]+}}OpFunctionEnd
+; CHECK-NOT: OpUnreachable
+
+declare void @llvm.trap() #0
+declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1
+declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1
+
+define spir_func void @trap_then_unreachable() {
+entry:
+ call void @llvm.trap()
+ unreachable
+}
+
+define spir_func void @trap_then_ret() {
+entry:
+ call void @llvm.trap()
+ ret void
+}
+
+define spir_func void @trap_then_lifetime_ret() {
+entry:
+ %buf = alloca i8, align 1
+ call void @llvm.lifetime.start.p0(i64 1, ptr %buf)
+ call void @llvm.trap()
+ call void @llvm.lifetime.end.p0(i64 1, ptr %buf)
+ ret void
+}
+
+attributes #0 = { cold noreturn nounwind }
+attributes #1 = { argmemonly nounwind willreturn }
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/ubsantrap-basic.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/ubsantrap-basic.ll
new file mode 100644
index 0000000000000..93f950f0b111b
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/ubsantrap-basic.ll
@@ -0,0 +1,25 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_abort %s -o - -filetype=obj | spirv-val %}
+
+;; llvm.ubsantrap(i8 N) lowers to OpAbortKHR with the i8 failure-kind argument
+;; zero-extended to a 32-bit Message operand.
+
+; CHECK-DAG: OpCapability AbortKHR
+; CHECK-DAG: OpExtension "SPV_KHR_abort"
+; CHECK-DAG: %[[#I32:]] = OpTypeInt 32 0
+; CHECK-DAG: %[[#MSG:]] = OpConstant %[[#I32]] 7
+
+; CHECK: OpAbortKHR %[[#I32]] %[[#MSG]]
+; CHECK-NOT: OpUnreachable
+; CHECK-NEXT: OpFunctionEnd
+
+define spir_func void @ubsantrap_simple() {
+entry:
+ call void @llvm.ubsantrap(i8 7)
+ unreachable
+}
+
+declare void @llvm.ubsantrap(i8) #0
+
+attributes #0 = { cold noreturn nounwind }
>From f4c0835bf50cade0193e8a2211f71f14d852c5db Mon Sep 17 00:00:00 2001
From: Victor Mustya <victor.mustya at intel.com>
Date: Thu, 30 Apr 2026 08:43:49 -0700
Subject: [PATCH 6/7] Fix formatting
---
llvm/lib/Target/SPIRV/SPIRVCommandLine.cpp | 3 +--
1 file changed, 1 insertion(+), 2 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVCommandLine.cpp b/llvm/lib/Target/SPIRV/SPIRVCommandLine.cpp
index 264fca1723442..a05d28470c793 100644
--- a/llvm/lib/Target/SPIRV/SPIRVCommandLine.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVCommandLine.cpp
@@ -177,8 +177,7 @@ static const StringMap<SPIRV::Extension::Extension> SPIRVExtensionMap = {
SPIRV::Extension::Extension::SPV_EXT_image_raw10_raw12},
{"SPV_INTEL_unstructured_loop_controls",
SPIRV::Extension::Extension::SPV_INTEL_unstructured_loop_controls},
- {"SPV_AMD_weak_linkage",
- SPIRV::Extension::Extension::SPV_AMD_weak_linkage},
+ {"SPV_AMD_weak_linkage", SPIRV::Extension::Extension::SPV_AMD_weak_linkage},
{"SPV_KHR_abort", SPIRV::Extension::Extension::SPV_KHR_abort}};
bool SPIRVExtensionsParser::parse(cl::Option &O, StringRef ArgName,
>From f65c7f7dd17cea8b2c80d9b95af1db8f62370190 Mon Sep 17 00:00:00 2001
From: Victor Mustya <victor.mustya at intel.com>
Date: Tue, 5 May 2026 07:23:41 -0700
Subject: [PATCH 7/7] Update llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
Co-authored-by: Marcos Maronas <mmaronas at amd.com>
---
llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index 0e7d790d4340f..c7bc25869fdd9 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -1744,7 +1744,7 @@ bool SPIRVInstructionSelector::selectPopCount(Register ResVReg,
MachineInstr &I,
unsigned Opcode) const {
// Vulkan restricts OpBitCount to 32-bit integers or vectors of 32-bit
- // integers unless VK_KHR_maintenance9 is enabled. Until VK_KHR_maintaince9
+ // integers unless VK_KHR_maintenance9 is enabled. Until VK_KHR_maintenance9
// is core we will not generate OpBitCount with any other types when
// targeting Vulkan.
if (!STI.getTargetTriple().isVulkanOS())
More information about the llvm-commits
mailing list