diff --git a/llvm/docs/SPIRVUsage.rst b/llvm/docs/SPIRVUsage.rst
index 2e67cf3eb99b0..0e633e7771c8f 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 44e31a1410523..5fef26d859d03 100644
--- a/llvm/include/llvm/IR/IntrinsicsSPIRV.td
+++ b/llvm/include/llvm/IR/IntrinsicsSPIRV.td
@@ -41,6 +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_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/SPIRVCommandLine.cpp b/llvm/lib/Target/SPIRV/SPIRVCommandLine.cpp
index 3d13cc01a3c18..a05d28470c793 100644
--- a/llvm/lib/Target/SPIRV/SPIRVCommandLine.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVCommandLine.cpp
@@ -177,8 +177,8 @@ 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,
                                   StringRef ArgValue, ExtensionSet &Vals) {
diff --git a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
index 180a0f45874e1..57147076402f7 100644
--- a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
@@ -1550,6 +1550,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());
@@ -2360,9 +2376,47 @@ Instruction *SPIRVEmitIntrinsics::visitAtomicCmpXchgInst(AtomicCmpXchgInst &I) {
   return NewI;
 }
 
+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 (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) {
-  IRBuilder<> B(I.getParent());
-  B.SetInsertPoint(&I);
+  const SPIRVSubtarget &ST = TM.getSubtarget<SPIRVSubtarget>(*I.getFunction());
+  if (precededByAbortIntrinsic(I, ST))
+    return &I;
+  IRBuilder<> B(&I);
   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 aee3a29c6e42b..c7bc25869fdd9 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -348,6 +348,8 @@ 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,
@@ -1375,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:
@@ -1739,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_maintenance9
+  // is core we will not generate OpBitCount with any other types when
   // targeting Vulkan.
   if (!STI.getTargetTriple().isVulkanOS())
     return selectUnOp(ResVReg, ResType, I, Opcode);
@@ -4760,6 +4764,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:
@@ -6326,6 +6332,86 @@ bool SPIRVInstructionSelector::selectAllocaArray(Register ResVReg,
   return true;
 }
 
+// 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 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: {
+      Register OperandReg = T->getOperand(1).getReg();
+      SPIRVTypeInst ElementT = GR.getSPIRVTypeForVReg(OperandReg);
+      Worklist.push_back(ElementT);
+    } break;
+    case SPIRV::OpTypeStruct:
+      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:
+      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))
+      .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 f32f3006c5e0e..c3c07f2efd6dd 100644
--- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
@@ -1967,6 +1967,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 10c9e20414692..916353ffdf381 100644
--- a/llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp
@@ -20,6 +20,7 @@
 
 #include "SPIRVPrepareFunctions.h"
 #include "SPIRV.h"
+#include "SPIRVBuiltins.h"
 #include "SPIRVSubtarget.h"
 #include "SPIRVTargetMachine.h"
 #include "SPIRVUtils.h"
@@ -34,6 +35,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>
 
@@ -44,6 +46,8 @@ namespace {
 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);
 
@@ -540,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());
@@ -605,6 +611,107 @@ SPIRVPrepareFunctionsImpl::removeAggregateTypesFromSignature(Function *F) {
   return NewF;
 }
 
+// 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();
+  // Quick reject: the mangled or unmangled name must contain the substring.
+  if (!Name.contains("__spirv_AbortKHR"))
+    return false;
+  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 (!isAbortKHRBuiltin(*F))
+    return false;
+
+  SmallVector<CallInst *> Calls;
+  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);
+  }
+
+  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) {
+    Instruction *Next = CI->getNextNode();
+    if (!Next || isa<UnreachableInst>(Next))
+      continue;
+    changeToUnreachable(Next);
+    Changed = true;
+  }
+  return Changed;
+}
+
 static std::string fixMultiOutputConstraintString(StringRef Constraints) {
   // We should only have one =r return for the made up ASM type.
   SmallVector<StringRef> Tmp;
@@ -719,7 +826,11 @@ bool SPIRVPrepareFunctionsImpl::runOnModule(Module &M) {
     Changed = true;
   }
 
+  Changed |= terminateBlocksAfterTrap(M, Intrinsic::trap);
+  Changed |= terminateBlocksAfterTrap(M, Intrinsic::ubsantrap);
+
   for (Function &F : M) {
+    Changed |= substituteAbortKHRCalls(&F);
     Changed |= substituteIntrinsicCalls(&F);
     Changed |= sortBlocks(F);
     Changed |= removeAggregateTypesFromCalls(&F);
diff --git a/llvm/lib/Target/SPIRV/SPIRVSymbolicOperands.td b/llvm/lib/Target/SPIRV/SPIRVSymbolicOperands.td
index cd52181bfc436..07b5d37f12580 100644
--- a/llvm/lib/Target/SPIRV/SPIRVSymbolicOperands.td
+++ b/llvm/lib/Target/SPIRV/SPIRVSymbolicOperands.td
@@ -399,6 +399,10 @@ defm SPV_ALTERA_arbitrary_precision_floating_point: ExtensionOperand<134, [EnvOp
 defm SPV_KHR_fma : ExtensionOperand<135, [EnvVulkan, EnvOpenCL]>;
 defm SPV_INTEL_masked_gather_scatter : ExtensionOperand<136, [EnvOpenCL]>;
 defm SPV_AMD_weak_linkage : ExtensionOperand<137, [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<138, [EnvOpenCL]>;
 
 //===----------------------------------------------------------------------===//
 // Multiclass used to define Capabilities enum values and at the same time
@@ -599,6 +603,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..9bea2744ac281
--- /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(<2 x i32>) #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..fd352bffd87cd
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-composite.ll
@@ -0,0 +1,52 @@
+; 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 @_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 @_Z16__spirv_AbortKHR3Msg(%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 @_Z16__spirv_AbortKHRDv4_j(<4 x i32> %v3)
+  unreachable
+}
+
+attributes #0 = { noreturn }
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
new file mode 100644
index 0000000000000..f5c2dcc34d0c5
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-invalid-message-type.ll
@@ -0,0 +1,26 @@
+; RUN: split-file %s %t
+; 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-bool.ll
+declare void @llvm.spv.abort.i1(i1) #0
+define void @abort_with_bool(i1 %b) {
+entry:
+  call void @llvm.spv.abort.i1(i1 %b)
+  unreachable
+}
+attributes #0 = { noreturn }
+
+;--- nested-bool.ll
+%B = type { i32, i1 }
+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.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-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-pointer.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/abort-pointer.ll
new file mode 100644
index 0000000000000..3693091759610
--- /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.p1(ptr addrspace(1)) #0
+
+define spir_kernel void @abort_with_pointer(ptr addrspace(1) %p) {
+entry:
+  call void @llvm.spv.abort.p1(ptr addrspace(1) %p)
+  unreachable
+}
+
+attributes #0 = { noreturn }
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
new file mode 100644
index 0000000000000..da1e20ea66877
--- /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(i32)
+
+define void @abort_with_int() {
+entry:
+  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..38669ae21646b
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_abort/ubsantrap-basic.ll
@@ -0,0 +1,31 @@
+; 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.ubsantrap is dropped silently (existing behavior).
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-NO-EXT
+
+;; 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
+
+; CHECK-NO-EXT-NOT: OpCapability AbortKHR
+; CHECK-NO-EXT-NOT: OpAbortKHR
+
+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 }