[SPIRV] Add Lifetime intrinsics/instructions #85391
Merged
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
|
@llvm/pr-subscribers-llvm-ir @llvm/pr-subscribers-mc Author: Vyacheslav Levytskyy (VyacheslavLevytskyy) ChangesThis PR:
Full diff: https://github.com/llvm/llvm-project/pull/85391.diff 17 Files Affected:
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..028b5df31e925e 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,13 @@ 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 +176,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..08e55d10272e54 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,27 @@ 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>
|
|
✅ With the latest revision this PR passed the C/C++ code formatter. |
michalpaszkowski
approved these changes
Mar 17, 2024
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This PR: