[llvm] 59f34e8 - [SPIRV] Add Lifetime intrinsics/instructions (#85391)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Mar 18 03:42:48 PDT 2024
Author: Vyacheslav Levytskyy
Date: 2024-03-18T11:42:44+01:00
New Revision: 59f34e8c2b5463dcf39fbafffeb30a84ef7b6887
URL: https://github.com/llvm/llvm-project/commit/59f34e8c2b5463dcf39fbafffeb30a84ef7b6887
DIFF: https://github.com/llvm/llvm-project/commit/59f34e8c2b5463dcf39fbafffeb30a84ef7b6887.diff
LOG: [SPIRV] Add Lifetime intrinsics/instructions (#85391)
This PR:
* adds Lifetime intrinsics/instructions
* fixes how the binary header is emitted (correct version and better
approximation of Bound)
* add validation into more test cases
Added:
llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll
Modified:
llvm/include/llvm/IR/IntrinsicsSPIRV.td
llvm/lib/MC/SPIRVObjectWriter.cpp
llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp
llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.h
llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
llvm/lib/Target/SPIRV/SPIRVMCInstLower.cpp
llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.h
llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp
llvm/lib/Target/SPIRV/SPIRVSubtarget.h
llvm/test/CodeGen/SPIRV/ComparePointers.ll
llvm/test/CodeGen/SPIRV/transcoding/AtomicCompareExchangeExplicit_cl20.ll
llvm/test/CodeGen/SPIRV/transcoding/builtin_vars.ll
llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_arithmetics.ll
llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_opt.ll
Removed:
################################################################################
diff --git a/llvm/include/llvm/IR/IntrinsicsSPIRV.td b/llvm/include/llvm/IR/IntrinsicsSPIRV.td
index 766b3b542a9e06..0eb09b1699aff4 100644
--- a/llvm/include/llvm/IR/IntrinsicsSPIRV.td
+++ b/llvm/include/llvm/IR/IntrinsicsSPIRV.td
@@ -40,6 +40,18 @@ let TargetPrefix = "spv" in {
def int_spv_assume : Intrinsic<[], [llvm_i1_ty]>;
def int_spv_expect : Intrinsic<[llvm_anyint_ty], [LLVMMatchType<0>, LLVMMatchType<0>]>;
+ // Memory Use Markers
+ def int_spv_lifetime_start : Intrinsic<[],
+ [llvm_i64_ty, llvm_anyptr_ty],
+ [IntrArgMemOnly, IntrWillReturn,
+ NoCapture<ArgIndex<1>>,
+ ImmArg<ArgIndex<0>>]>;
+ def int_spv_lifetime_end : Intrinsic<[],
+ [llvm_i64_ty, llvm_anyptr_ty],
+ [IntrArgMemOnly, IntrWillReturn,
+ NoCapture<ArgIndex<1>>,
+ ImmArg<ArgIndex<0>>]>;
+
// The following intrinsic(s) are mirrored from IntrinsicsDirectX.td for HLSL support.
def int_spv_thread_id : Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrWillReturn]>;
def int_spv_create_handle : ClangBuiltin<"__builtin_hlsl_create_handle">,
diff --git a/llvm/lib/MC/SPIRVObjectWriter.cpp b/llvm/lib/MC/SPIRVObjectWriter.cpp
index 39856e96e9be51..d72d6e07f2e6fd 100644
--- a/llvm/lib/MC/SPIRVObjectWriter.cpp
+++ b/llvm/lib/MC/SPIRVObjectWriter.cpp
@@ -43,18 +43,14 @@ class SPIRVObjectWriter : public MCObjectWriter {
void SPIRVObjectWriter::writeHeader(const MCAssembler &Asm) {
constexpr uint32_t MagicNumber = 0x07230203;
-
- // TODO: set the version on a min-necessary basis (just like the translator
- // does) requires some refactoring of MCAssembler::VersionInfoType.
- constexpr uint32_t Major = 1;
- constexpr uint32_t Minor = 0;
- constexpr uint32_t VersionNumber = 0 | (Major << 16) | (Minor << 8);
- // TODO: check if we could use anything other than 0 (spec allows).
constexpr uint32_t GeneratorMagicNumber = 0;
- // TODO: do not hardcode this as well.
- constexpr uint32_t Bound = 900;
constexpr uint32_t Schema = 0;
+ // Construct SPIR-V version and Bound
+ const MCAssembler::VersionInfoType &VIT = Asm.getVersionInfo();
+ uint32_t VersionNumber = 0 | (VIT.Major << 16) | (VIT.Minor << 8);
+ uint32_t Bound = VIT.Update;
+
W.write<uint32_t>(MagicNumber);
W.write<uint32_t>(VersionNumber);
W.write<uint32_t>(GeneratorMagicNumber);
diff --git a/llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp b/llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp
index 1fbf3c3e11aedc..30c67d3fde6338 100644
--- a/llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp
@@ -29,7 +29,9 @@
#include "llvm/CodeGen/MachineModuleInfo.h"
#include "llvm/CodeGen/TargetLoweringObjectFileImpl.h"
#include "llvm/MC/MCAsmInfo.h"
+#include "llvm/MC/MCAssembler.h"
#include "llvm/MC/MCInst.h"
+#include "llvm/MC/MCObjectStreamer.h"
#include "llvm/MC/MCStreamer.h"
#include "llvm/MC/MCSymbol.h"
#include "llvm/MC/TargetRegistry.h"
@@ -101,6 +103,21 @@ void SPIRVAsmPrinter::emitEndOfAsmFile(Module &M) {
if (ModuleSectionsEmitted == false) {
outputModuleSections();
ModuleSectionsEmitted = true;
+ } else {
+ ST = static_cast<const SPIRVTargetMachine &>(TM).getSubtargetImpl();
+ uint32_t DecSPIRVVersion = ST->getSPIRVVersion();
+ uint32_t Major = DecSPIRVVersion / 10;
+ uint32_t Minor = DecSPIRVVersion - Major * 10;
+ // TODO: calculate Bound more carefully from maximum used register number,
+ // accounting for generated OpLabels and other related instructions if
+ // needed.
+ unsigned Bound = 2 * (ST->getBound() + 1);
+ bool FlagToRestore = OutStreamer->getUseAssemblerInfoForParsing();
+ OutStreamer->setUseAssemblerInfoForParsing(true);
+ if (MCAssembler *Asm = OutStreamer->getAssemblerPtr())
+ Asm->setBuildVersion(static_cast<MachO::PlatformType>(0), Major, Minor,
+ Bound, VersionTuple(Major, Minor, 0, Bound));
+ OutStreamer->setUseAssemblerInfoForParsing(FlagToRestore);
}
}
@@ -507,6 +524,13 @@ void SPIRVAsmPrinter::outputAnnotations(const Module &M) {
report_fatal_error("Unsupported value in llvm.global.annotations");
Function *Func = cast<Function>(AnnotatedVar);
Register Reg = MAI->getFuncReg(Func);
+ if (!Reg.isValid()) {
+ std::string DiagMsg;
+ raw_string_ostream OS(DiagMsg);
+ AnnotatedVar->print(OS);
+ DiagMsg = "Unknown function in llvm.global.annotations: " + DiagMsg;
+ report_fatal_error(DiagMsg.c_str());
+ }
// The second field contains a pointer to a global annotation string.
GlobalVariable *GV =
diff --git a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
index bda9c57e534c3a..42f8397a3023b1 100644
--- a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
@@ -24,7 +24,7 @@
using namespace llvm;
SPIRVGlobalRegistry::SPIRVGlobalRegistry(unsigned PointerSize)
- : PointerSize(PointerSize) {}
+ : PointerSize(PointerSize), Bound(0) {}
SPIRVType *SPIRVGlobalRegistry::assignIntTypeToVReg(unsigned BitWidth,
Register VReg,
@@ -896,6 +896,15 @@ bool SPIRVGlobalRegistry::isScalarOrVectorSigned(const SPIRVType *Type) const {
return IntType && IntType->getOperand(2).getImm() != 0;
}
+unsigned SPIRVGlobalRegistry::getPointeeTypeOp(Register PtrReg) {
+ SPIRVType *PtrType = getSPIRVTypeForVReg(PtrReg);
+ SPIRVType *ElemType =
+ PtrType && PtrType->getOpcode() == SPIRV::OpTypePointer
+ ? getSPIRVTypeForVReg(PtrType->getOperand(2).getReg())
+ : nullptr;
+ return ElemType ? ElemType->getOpcode() : 0;
+}
+
bool SPIRVGlobalRegistry::isBitcastCompatible(const SPIRVType *Type1,
const SPIRVType *Type2) const {
if (!Type1 || !Type2)
diff --git a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.h b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.h
index 25d82ebf9bc79b..da480b22a525f2 100644
--- a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.h
+++ b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.h
@@ -56,6 +56,9 @@ class SPIRVGlobalRegistry {
// Number of bits pointers and size_t integers require.
const unsigned PointerSize;
+ // Holds the maximum ID we have in the module.
+ unsigned Bound;
+
// Add a new OpTypeXXX instruction without checking for duplicates.
SPIRVType *createSPIRVType(const Type *Type, MachineIRBuilder &MIRBuilder,
SPIRV::AccessQualifier::AccessQualifier AQ =
@@ -108,6 +111,9 @@ class SPIRVGlobalRegistry {
DT.buildDepsGraph(Graph, MMI);
}
+ void setBound(unsigned V) { Bound = V; }
+ unsigned getBound() { return Bound; }
+
// Map a machine operand that represents a use of a function via function
// pointer to a machine operand that represents the function definition.
// Return either the register or invalid value, because we have no context for
@@ -166,6 +172,9 @@ class SPIRVGlobalRegistry {
return Res->second;
}
+ // Return a pointee's type op code, or 0 otherwise.
+ unsigned getPointeeTypeOp(Register PtrReg);
+
// Either generate a new OpTypeXXX instruction or return an existing one
// corresponding to the given string containing the name of the builtin type.
// Return nullptr if unable to recognize SPIRV type name from `TypeStr`.
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index fd19b7412c4c9c..0fef19c2d53419 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -1567,7 +1567,8 @@ bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
MachineBasicBlock &BB = *I.getParent();
- switch (cast<GIntrinsic>(I).getIntrinsicID()) {
+ Intrinsic::ID IID = cast<GIntrinsic>(I).getIntrinsicID();
+ switch (IID) {
case Intrinsic::spv_load:
return selectLoad(ResVReg, ResType, I);
case Intrinsic::spv_store:
@@ -1661,8 +1662,25 @@ bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
break;
case Intrinsic::spv_thread_id:
return selectSpvThreadId(ResVReg, ResType, I);
- default:
- llvm_unreachable("Intrinsic selection not implemented");
+ case Intrinsic::spv_lifetime_start:
+ case Intrinsic::spv_lifetime_end: {
+ unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
+ : SPIRV::OpLifetimeStop;
+ int64_t Size = I.getOperand(I.getNumExplicitDefs() + 1).getImm();
+ Register PtrReg = I.getOperand(I.getNumExplicitDefs() + 2).getReg();
+ unsigned PonteeOpType = GR.getPointeeTypeOp(PtrReg);
+ bool IsNonvoidPtr = PonteeOpType != 0 && PonteeOpType != SPIRV::OpTypeVoid;
+ if (Size == -1 || IsNonvoidPtr)
+ Size = 0;
+ BuildMI(BB, I, I.getDebugLoc(), TII.get(Op)).addUse(PtrReg).addImm(Size);
+ } break;
+ default: {
+ std::string DiagMsg;
+ raw_string_ostream OS(DiagMsg);
+ I.print(OS);
+ DiagMsg = "Intrinsic selection not implemented: " + DiagMsg;
+ report_fatal_error(DiagMsg.c_str(), false);
+ }
}
return true;
}
diff --git a/llvm/lib/Target/SPIRV/SPIRVMCInstLower.cpp b/llvm/lib/Target/SPIRV/SPIRVMCInstLower.cpp
index 8c6649bf628265..afa550d6dd424e 100644
--- a/llvm/lib/Target/SPIRV/SPIRVMCInstLower.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVMCInstLower.cpp
@@ -34,7 +34,13 @@ void SPIRVMCInstLower::lower(const MachineInstr *MI, MCInst &OutMI,
llvm_unreachable("unknown operand type");
case MachineOperand::MO_GlobalAddress: {
Register FuncReg = MAI->getFuncReg(dyn_cast<Function>(MO.getGlobal()));
- assert(FuncReg.isValid() && "Cannot find function Id");
+ if (!FuncReg.isValid()) {
+ std::string DiagMsg;
+ raw_string_ostream OS(DiagMsg);
+ MI->print(OS);
+ DiagMsg = "Unknown function in:" + DiagMsg;
+ report_fatal_error(DiagMsg.c_str());
+ }
MCOp = MCOperand::createReg(FuncReg);
break;
}
diff --git a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
index 2b4cb5ccc7b1eb..00d0cbd763736d 100644
--- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
@@ -1309,5 +1309,8 @@ bool SPIRVModuleAnalysis::runOnModule(Module &M) {
if (MAI.MS[SPIRV::MB_EntryPoints].empty())
MAI.Reqs.addCapability(SPIRV::Capability::Linkage);
+ // Set maximum ID used.
+ GR->setBound(MAI.MaxID);
+
return false;
}
diff --git a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.h b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.h
index 708384fc55f525..6e86eed30c5dc1 100644
--- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.h
+++ b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.h
@@ -163,8 +163,8 @@ struct ModuleAnalysisInfo {
Register getFuncReg(const Function *F) {
assert(F && "Function is null");
auto FuncPtrRegPair = FuncMap.find(F);
- assert(FuncPtrRegPair != FuncMap.end() && "Cannot find function ID");
- return FuncPtrRegPair->second;
+ return FuncPtrRegPair == FuncMap.end() ? Register(0)
+ : FuncPtrRegPair->second;
}
Register getExtInstSetReg(unsigned SetNum) { return ExtInstSetMap[SetNum]; }
InstrList &getMSInstrs(unsigned MSType) { return MS[MSType]; }
diff --git a/llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp b/llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp
index c376497469ce33..a8a0577f60564c 100644
--- a/llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp
@@ -263,6 +263,21 @@ static void lowerExpectAssume(IntrinsicInst *II) {
return;
}
+static bool toSpvOverloadedIntrinsic(IntrinsicInst *II, Intrinsic::ID NewID,
+ ArrayRef<unsigned> OpNos) {
+ Function *F = nullptr;
+ if (OpNos.empty()) {
+ F = Intrinsic::getDeclaration(II->getModule(), NewID);
+ } else {
+ SmallVector<Type *, 4> Tys;
+ for (unsigned OpNo : OpNos)
+ Tys.push_back(II->getOperand(OpNo)->getType());
+ F = Intrinsic::getDeclaration(II->getModule(), NewID, Tys);
+ }
+ II->setCalledFunction(F);
+ return true;
+}
+
static void lowerUMulWithOverflow(IntrinsicInst *UMulIntrinsic) {
// Get a separate function - otherwise, we'd have to rework the CFG of the
// current one. Then simply replace the intrinsic uses with a call to the new
@@ -290,22 +305,35 @@ bool SPIRVPrepareFunctions::substituteIntrinsicCalls(Function *F) {
if (!CF || !CF->isIntrinsic())
continue;
auto *II = cast<IntrinsicInst>(Call);
- if (II->getIntrinsicID() == Intrinsic::memset ||
- II->getIntrinsicID() == Intrinsic::bswap)
+ switch (II->getIntrinsicID()) {
+ case Intrinsic::memset:
+ case Intrinsic::bswap:
Changed |= lowerIntrinsicToFunction(II);
- else if (II->getIntrinsicID() == Intrinsic::fshl ||
- II->getIntrinsicID() == Intrinsic::fshr) {
+ break;
+ case Intrinsic::fshl:
+ case Intrinsic::fshr:
lowerFunnelShifts(II);
Changed = true;
- } else if (II->getIntrinsicID() == Intrinsic::umul_with_overflow) {
+ break;
+ case Intrinsic::umul_with_overflow:
lowerUMulWithOverflow(II);
Changed = true;
- } else if (II->getIntrinsicID() == Intrinsic::assume ||
- II->getIntrinsicID() == Intrinsic::expect) {
+ break;
+ case Intrinsic::assume:
+ case Intrinsic::expect: {
const SPIRVSubtarget &STI = TM.getSubtarget<SPIRVSubtarget>(*F);
if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_expect_assume))
lowerExpectAssume(II);
Changed = true;
+ } break;
+ case Intrinsic::lifetime_start:
+ Changed |= toSpvOverloadedIntrinsic(
+ II, Intrinsic::SPVIntrinsics::spv_lifetime_start, {1});
+ break;
+ case Intrinsic::lifetime_end:
+ Changed |= toSpvOverloadedIntrinsic(
+ II, Intrinsic::SPVIntrinsics::spv_lifetime_end, {1});
+ break;
}
}
}
diff --git a/llvm/lib/Target/SPIRV/SPIRVSubtarget.h b/llvm/lib/Target/SPIRV/SPIRVSubtarget.h
index 62524ebfc9bf8c..3b486226a93931 100644
--- a/llvm/lib/Target/SPIRV/SPIRVSubtarget.h
+++ b/llvm/lib/Target/SPIRV/SPIRVSubtarget.h
@@ -71,6 +71,7 @@ class SPIRVSubtarget : public SPIRVGenSubtargetInfo {
// The definition of this function is auto generated by tblgen.
void ParseSubtargetFeatures(StringRef CPU, StringRef TuneCPU, StringRef FS);
unsigned getPointerSize() const { return PointerSize; }
+ unsigned getBound() const { return GR->getBound(); }
bool canDirectlyComparePointers() const;
// TODO: this environment is not implemented in Triple, we need to decide
// how to standardize its support. For now, let's assume SPIR-V with physical
diff --git a/llvm/test/CodeGen/SPIRV/ComparePointers.ll b/llvm/test/CodeGen/SPIRV/ComparePointers.ll
index 9be05944789b6f..6777fc38024b32 100644
--- a/llvm/test/CodeGen/SPIRV/ComparePointers.ll
+++ b/llvm/test/CodeGen/SPIRV/ComparePointers.ll
@@ -1,5 +1,5 @@
; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --mattr=+spirv1.3 %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
;; kernel void test(int global *in, int global *in2) {
;; if (!in)
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll
new file mode 100644
index 00000000000000..710a1581f760ca
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll
@@ -0,0 +1,25 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK: OpFunction
+; CHECK: %[[FooArg:.*]] = OpVariable
+; CHECK: OpLifetimeStart %[[FooArg]], 0
+; CHECK: OpCopyMemorySized
+; CHECK: OpBitcast
+; CHECK: OpInBoundsPtrAccessChain
+; CHECK: OpLifetimeStop %[[FooArg]], 0
+
+%tprange = type { %tparray }
+%tparray = type { [2 x i64] }
+
+define spir_func void @foo(ptr noundef byval(%tprange) align 8 %_arg_UserRange) {
+ %RoundedRangeKernel = alloca %tprange, align 8
+ call void @llvm.lifetime.start.p0(i64 72, ptr nonnull %RoundedRangeKernel) #7
+ call void @llvm.memcpy.p0.p0.i64(ptr align 8 %RoundedRangeKernel, ptr align 8 %_arg_UserRange, i64 16, i1 false)
+ %KernelFunc = getelementptr inbounds i8, ptr %RoundedRangeKernel, i64 16
+ call void @llvm.lifetime.end.p0(i64 72, ptr nonnull %RoundedRangeKernel) #7
+ ret void
+}
+
+declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture)
+declare void @llvm.memcpy.p0.p0.i64(ptr noalias nocapture writeonly, ptr noalias nocapture readonly, i64, i1 immarg)
+declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture)
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/AtomicCompareExchangeExplicit_cl20.ll b/llvm/test/CodeGen/SPIRV/transcoding/AtomicCompareExchangeExplicit_cl20.ll
index 55cfcea999d84b..e0c47798cc6d09 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/AtomicCompareExchangeExplicit_cl20.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/AtomicCompareExchangeExplicit_cl20.ll
@@ -1,5 +1,5 @@
-; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --mattr=+spirv1.3 %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown --mattr=+spirv1.3 %s -o - -filetype=obj | spirv-val %}
;; __kernel void testAtomicCompareExchangeExplicit_cl20(
;; volatile global atomic_int* object,
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars.ll b/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars.ll
index f18f27a6de51d4..50748931635656 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars.ll
@@ -1,5 +1,5 @@
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
; CHECK-SPIRV: OpDecorate %[[#Id:]] BuiltIn GlobalLinearId
; CHECK-SPIRV: %[[#Id:]] = OpVariable %[[#]]
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_arithmetics.ll b/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_arithmetics.ll
index d39ca3c39383c0..d0c4dff43121cc 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_arithmetics.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_arithmetics.ll
@@ -1,5 +1,5 @@
; RUN: llc -O0 -mtriple=spirv64-unknown-linux %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
;; The IR was generated from the following source:
;; #include <CL/sycl.hpp>
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_opt.ll b/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_opt.ll
index 03456aef6b6b2e..3885f070231442 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_opt.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_opt.ll
@@ -1,5 +1,5 @@
; RUN: llc -O0 -mtriple=spirv64-unknown-linux %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
;; The IR was generated from the following source:
;; #include <CL/sycl.hpp>
More information about the llvm-commits
mailing list