Skip to content

Commit 5d28c93

Browse files
asudarsaasudarsa
authored
[llvm-spirv] Cherry pick Khronos changes to expand collection of entry point interfaces (PR #1334) (#10623)
This PR pulls in the following PR from upstream Khronos SPIRV-LLVM-Translator repo: KhronosGroup/SPIRV-LLVM-Translator#1334 ` This is a patch to expand the collection of entry point interfaces. In SPIR-V 1.4 and later OpEntryPoint must list all global variables in the interface. ` In addition, a couple of minor changes have been added to sync with latest code. This patch addresses #9958 Updated the following tests to sync with upstream as well: llvm-spirv/test/extensions/INTEL/SPV_INTEL_inline_assembly/inline_asm_clobbers.cl llvm-spirv/test/extensions/INTEL/SPV_INTEL_inline_assembly/inline_asm_constraints.cl Thanks --------- Signed-off-by: Arvind Sudarsanam <arvind.sudarsanam@intel.com>
1 parent dd00ff6 commit 5d28c93

File tree

15 files changed

+107
-62
lines changed

15 files changed

+107
-62
lines changed

‎llvm-spirv/lib/SPIRV/SPIRVWriter.cpp

Lines changed: 18 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -874,9 +874,7 @@ SPIRVFunction *LLVMToSPIRVBase::transFunctionDecl(Function *F) {
874874

875875
BM->setName(BF, F->getName().str());
876876
}
877-
if (isKernel(F))
878-
BM->addEntryPoint(ExecutionModelKernel, BF->getId());
879-
else if (F->getLinkage() != GlobalValue::InternalLinkage)
877+
if (!isKernel(F) && F->getLinkage() != GlobalValue::InternalLinkage)
880878
BF->setLinkageType(transLinkageType(F));
881879

882880
// Translate OpenCL/SYCL buffer_location metadata if it's attached to the
@@ -4898,12 +4896,15 @@ bool LLVMToSPIRVBase::isAnyFunctionReachableFromFunction(
48984896
return false;
48994897
}
49004898

4901-
void LLVMToSPIRVBase::collectInputOutputVariables(SPIRVFunction *SF,
4902-
Function *F) {
4899+
std::vector<SPIRVId>
4900+
LLVMToSPIRVBase::collectEntryPointInterfaces(SPIRVFunction *SF, Function *F) {
4901+
std::vector<SPIRVId> Interface;
49034902
for (auto &GV : M->globals()) {
49044903
const auto AS = GV.getAddressSpace();
4905-
if (AS != SPIRAS_Input && AS != SPIRAS_Output)
4906-
continue;
4904+
SPIRVModule *BM = SF->getModule();
4905+
if (!BM->isAllowedToUseVersion(VersionNumber::SPIRV_1_4))
4906+
if (AS != SPIRAS_Input && AS != SPIRAS_Output)
4907+
continue;
49074908

49084909
std::unordered_set<const Function *> Funcs;
49094910

@@ -4915,9 +4916,14 @@ void LLVMToSPIRVBase::collectInputOutputVariables(SPIRVFunction *SF,
49154916
}
49164917

49174918
if (isAnyFunctionReachableFromFunction(F, Funcs)) {
4918-
SF->addVariable(ValueMap[&GV]);
4919+
SPIRVWord ModuleVersion = static_cast<SPIRVWord>(BM->getSPIRVVersion());
4920+
if (AS != SPIRAS_Input && AS != SPIRAS_Output &&
4921+
ModuleVersion < static_cast<SPIRVWord>(VersionNumber::SPIRV_1_4))
4922+
BM->setMinSPIRVVersion(VersionNumber::SPIRV_1_4);
4923+
Interface.push_back(ValueMap[&GV]->getId());
49194924
}
49204925
}
4926+
return Interface;
49214927
}
49224928

49234929
void LLVMToSPIRVBase::mutateFuncArgType(
@@ -5118,10 +5124,10 @@ void LLVMToSPIRVBase::transFunction(Function *I) {
51185124
joinFPContract(I, FPContract::ENABLED);
51195125
fpContractUpdateRecursive(I, getFPContract(I));
51205126

5121-
bool IsKernelEntryPoint = isKernel(I);
5122-
5123-
if (IsKernelEntryPoint) {
5124-
collectInputOutputVariables(BF, I);
5127+
if (isKernel(I)) {
5128+
auto Interface = collectEntryPointInterfaces(BF, I);
5129+
BM->addEntryPoint(ExecutionModelKernel, BF->getId(), I->getName().str(),
5130+
Interface);
51255131
}
51265132
}
51275133

‎llvm-spirv/lib/SPIRV/SPIRVWriter.h

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -248,6 +248,8 @@ class LLVMToSPIRVBase : protected BuiltinCallHelper {
248248
const Function *FS,
249249
const std::unordered_set<const Function *> Funcs) const;
250250
void collectInputOutputVariables(SPIRVFunction *SF, Function *F);
251+
std::vector<SPIRVId> collectEntryPointInterfaces(SPIRVFunction *BF,
252+
Function *F);
251253
};
252254

253255
class LLVMToSPIRVPass : public PassInfoMixin<LLVMToSPIRVPass> {

‎llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -268,7 +268,6 @@ class SPIRVDecorateLinkageAttr : public SPIRVDecorate {
268268
#ifdef _SPIRV_SUPPORT_TEXT_FMT
269269
if (SPIRVUseTextFormat) {
270270
Encoder << getString(Literals.cbegin(), Literals.cend() - 1);
271-
Encoder.OS << " ";
272271
Encoder << (SPIRVLinkageTypeKind)Literals.back();
273272
} else
274273
#endif

‎llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.cpp

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -621,9 +621,11 @@ void SPIRVEntryPoint::encode(spv_ostream &O) const {
621621
}
622622

623623
void SPIRVEntryPoint::decode(std::istream &I) {
624-
getDecoder(I) >> ExecModel >> Target >> Name >> Variables;
624+
getDecoder(I) >> ExecModel >> Target >> Name;
625+
Variables.resize(WordCount - FixedWC - getSizeInWords(Name) + 1);
626+
getDecoder(I) >> Variables;
625627
Module->setName(getOrCreateTarget(), Name);
626-
Module->addEntryPoint(ExecModel, Target);
628+
Module->addEntryPoint(ExecModel, Target, Name, Variables);
627629
}
628630

629631
void SPIRVExecutionMode::encode(spv_ostream &O) const {

‎llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.h

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -533,6 +533,7 @@ template <Op OC> class SPIRVAnnotation : public SPIRVAnnotationGeneric {
533533

534534
class SPIRVEntryPoint : public SPIRVAnnotation<OpEntryPoint> {
535535
public:
536+
static const SPIRVWord FixedWC = 4;
536537
SPIRVEntryPoint(SPIRVModule *TheModule, SPIRVExecutionModelKind,
537538
SPIRVId TheId, const std::string &TheName,
538539
std::vector<SPIRVId> Variables);

‎llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.cpp

Lines changed: 12 additions & 26 deletions
Original file line numberDiff line numberDiff line change
@@ -128,20 +128,6 @@ class SPIRVModuleImpl : public SPIRVModule {
128128
getValueTypes(const std::vector<SPIRVId> &) const override;
129129
SPIRVMemoryModelKind getMemoryModel() const override { return MemoryModel; }
130130
SPIRVConstant *getLiteralAsConstant(unsigned Literal) override;
131-
unsigned getNumEntryPoints(SPIRVExecutionModelKind EM) const override {
132-
auto Loc = EntryPointVec.find(EM);
133-
if (Loc == EntryPointVec.end())
134-
return 0;
135-
return Loc->second.size();
136-
}
137-
SPIRVFunction *getEntryPoint(SPIRVExecutionModelKind EM,
138-
unsigned I) const override {
139-
auto Loc = EntryPointVec.find(EM);
140-
if (Loc == EntryPointVec.end())
141-
return nullptr;
142-
assert(I < Loc->second.size());
143-
return get<SPIRVFunction>(Loc->second[I]);
144-
}
145131
unsigned getNumFunctions() const override { return FuncVec.size(); }
146132
unsigned getNumVariables() const override { return VariableVec.size(); }
147133
SourceLanguage getSourceLanguage(SPIRVWord *Ver = nullptr) const override {
@@ -225,8 +211,9 @@ class SPIRVModuleImpl : public SPIRVModule {
225211
SPIRVGroupMemberDecorate *
226212
addGroupMemberDecorate(SPIRVDecorationGroup *Group,
227213
const std::vector<SPIRVEntry *> &Targets) override;
228-
void addEntryPoint(SPIRVExecutionModelKind ExecModel,
229-
SPIRVId EntryPoint) override;
214+
void addEntryPoint(SPIRVExecutionModelKind ExecModel, SPIRVId EntryPoint,
215+
const std::string &Name,
216+
const std::vector<SPIRVId> &Variables) override;
230217
SPIRVForward *addForward(SPIRVType *Ty) override;
231218
SPIRVForward *addForward(SPIRVId, SPIRVType *Ty) override;
232219
SPIRVFunction *addFunction(SPIRVFunction *) override;
@@ -508,11 +495,11 @@ class SPIRVModuleImpl : public SPIRVModule {
508495
typedef std::vector<SPIRVGroupDecorateGeneric *> SPIRVGroupDecVec;
509496
typedef std::vector<SPIRVAsmTargetINTEL *> SPIRVAsmTargetVector;
510497
typedef std::vector<SPIRVAsmINTEL *> SPIRVAsmVector;
498+
typedef std::vector<SPIRVEntryPoint *> SPIRVEntryPointVec;
511499
typedef std::map<SPIRVId, SPIRVExtInstSetKind> SPIRVIdToInstructionSetMap;
512500
std::map<SPIRVExtInstSetKind, SPIRVId> ExtInstSetIds;
513501
typedef std::map<SPIRVId, SPIRVExtInstSetKind> SPIRVIdToBuiltinSetMap;
514502
typedef std::map<SPIRVExecutionModelKind, SPIRVIdSet> SPIRVExecModelIdSetMap;
515-
typedef std::map<SPIRVExecutionModelKind, SPIRVIdVec> SPIRVExecModelIdVecMap;
516503
typedef std::unordered_map<std::string, SPIRVString *> SPIRVStringMap;
517504
typedef std::map<SPIRVTypeStruct *, std::vector<std::pair<unsigned, SPIRVId>>>
518505
SPIRVUnknownStructFieldMap;
@@ -540,7 +527,7 @@ class SPIRVModuleImpl : public SPIRVModule {
540527
SPIRVAsmTargetVector AsmTargetVec;
541528
SPIRVAsmVector AsmVec;
542529
SPIRVExecModelIdSetMap EntryPointSet;
543-
SPIRVExecModelIdVecMap EntryPointVec;
530+
SPIRVEntryPointVec EntryPointVec;
544531
SPIRVStringMap StrMap;
545532
SPIRVCapMap CapMap;
546533
SPIRVUnknownStructFieldMap UnknownStructFieldMap;
@@ -1086,11 +1073,14 @@ SPIRVModuleImpl::addDecorate(SPIRVDecorateGeneric *Dec) {
10861073
}
10871074

10881075
void SPIRVModuleImpl::addEntryPoint(SPIRVExecutionModelKind ExecModel,
1089-
SPIRVId EntryPoint) {
1076+
SPIRVId EntryPoint, const std::string &Name,
1077+
const std::vector<SPIRVId> &Variables) {
10901078
assert(isValid(ExecModel) && "Invalid execution model");
10911079
assert(EntryPoint != SPIRVID_INVALID && "Invalid entry point");
1080+
auto *EP =
1081+
add(new SPIRVEntryPoint(this, ExecModel, EntryPoint, Name, Variables));
1082+
EntryPointVec.push_back(EP);
10921083
EntryPointSet[ExecModel].insert(EntryPoint);
1093-
EntryPointVec[ExecModel].push_back(EntryPoint);
10941084
addCapabilities(SPIRV::getCapability(ExecModel));
10951085
}
10961086

@@ -1937,14 +1927,10 @@ spv_ostream &operator<<(spv_ostream &O, SPIRVModule &M) {
19371927

19381928
O << SPIRVMemoryModel(&M);
19391929

1940-
for (auto &I : MI.EntryPointVec)
1941-
for (auto &II : I.second)
1942-
O << SPIRVEntryPoint(&M, I.first, II, M.get<SPIRVFunction>(II)->getName(),
1943-
M.get<SPIRVFunction>(II)->getVariables());
1930+
O << MI.EntryPointVec;
19441931

19451932
for (auto &I : MI.EntryPointVec)
1946-
for (auto &II : I.second)
1947-
MI.get<SPIRVFunction>(II)->encodeExecutionModes(O);
1933+
MI.get<SPIRVFunction>(I->getTargetId())->encodeExecutionModes(O);
19481934

19491935
O << MI.StringVec;
19501936

‎llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.h

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -133,14 +133,11 @@ class SPIRVModule {
133133
virtual const SPIRVCapMap &getCapability() const = 0;
134134
virtual bool hasCapability(SPIRVCapabilityKind) const = 0;
135135
virtual SPIRVExtInstSetKind getBuiltinSet(SPIRVId) const = 0;
136-
virtual SPIRVFunction *getEntryPoint(SPIRVExecutionModelKind,
137-
unsigned) const = 0;
138136
virtual std::set<std::string> &getExtension() = 0;
139137
virtual SPIRVFunction *getFunction(unsigned) const = 0;
140138
virtual SPIRVVariable *getVariable(unsigned) const = 0;
141139
virtual SPIRVMemoryModelKind getMemoryModel() const = 0;
142140
virtual unsigned getNumFunctions() const = 0;
143-
virtual unsigned getNumEntryPoints(SPIRVExecutionModelKind) const = 0;
144141
virtual unsigned getNumVariables() const = 0;
145142
virtual SourceLanguage getSourceLanguage(SPIRVWord *) const = 0;
146143
virtual std::set<std::string> &getSourceExtension() = 0;
@@ -223,7 +220,9 @@ class SPIRVModule {
223220
const std::vector<SPIRVEntry *> &Targets) = 0;
224221
virtual SPIRVGroupDecorateGeneric *
225222
addGroupDecorateGeneric(SPIRVGroupDecorateGeneric *GDec) = 0;
226-
virtual void addEntryPoint(SPIRVExecutionModelKind, SPIRVId) = 0;
223+
virtual void addEntryPoint(SPIRVExecutionModelKind, SPIRVId,
224+
const std::string &,
225+
const std::vector<SPIRVId> &) = 0;
227226
virtual SPIRVForward *addForward(SPIRVType *Ty) = 0;
228227
virtual SPIRVForward *addForward(SPIRVId, SPIRVType *Ty) = 0;
229228
virtual SPIRVFunction *addFunction(SPIRVFunction *) = 0;

‎llvm-spirv/lib/SPIRV/libSPIRV/SPIRVStream.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -180,6 +180,7 @@ const SPIRVEncoder &operator<<(const SPIRVEncoder &O, const std::string &Str) {
180180
#ifdef _SPIRV_SUPPORT_TEXT_FMT
181181
if (SPIRVUseTextFormat) {
182182
writeQuotedString(O.OS, Str);
183+
O.OS << " ";
183184
return O;
184185
}
185186
#endif

‎llvm-spirv/test/ExecutionMode.ll

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,9 +1,6 @@
11
; RUN: llvm-as < %s | llvm-spirv -spirv-text -o %t
22
; RUN: FileCheck < %t %s
33

4-
; check for magic number followed by version 1.1
5-
; CHECK: 119734787 65792
6-
74
; CHECK-DAG: TypeVoid [[VOID:[0-9]+]]
85

96
; CHECK-DAG: EntryPoint 6 [[WORKER:[0-9]+]] "worker"

‎llvm-spirv/test/copy_object.spt

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -5,7 +5,7 @@
55
2 Capability Int64
66
2 Capability Int8
77
3 MemoryModel 2 2
8-
8 EntryPoint 6 1 "copy_object"
8+
6 EntryPoint 6 1 "copy_object"
99
3 Source 3 102000
1010
3 Name 2 "in"
1111
4 Decorate 3 BuiltIn 28

‎llvm-spirv/test/entry-point-interfaces.ll

Lines changed: 52 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,52 @@
1+
; RUN: llvm-as %s -o %t.bc
2+
3+
; RUN: llvm-spirv %t.bc -o %t.spv
4+
; RUN: spirv-val --target-env spv1.4 %t.spv
5+
; RUN: llvm-spirv -to-text %t.spv -o %t.from.spv.spt
6+
; RUN: FileCheck < %t.from.spv.spt %s --check-prefix=CHECK-SPIRV
7+
8+
; RUN: llvm-spirv -spirv-text %t.bc -o %t.from.bc.spt
9+
; RUN: FileCheck < %t.from.bc.spt %s --check-prefix=CHECK-SPIRV
10+
11+
; CHECK-SPIRV: 7 EntryPoint 6 [[#]] "test" [[#Interface1:]] [[#Interface2:]]
12+
; CHECK-SPIRV: TypeInt [[#TypeInt:]] 32 0
13+
; CHECK-SPIRV: Constant [[#TypeInt]] [[#Constant1:]] 1
14+
; CHECK-SPIRV: Constant [[#TypeInt]] [[#Constant2:]] 3
15+
; CHECK-SPIRV: Variable [[#]] [[#Interface1]] 0 [[#Constant1]]
16+
; CHECK-SPIRV: Variable [[#]] [[#Interface2]] 0 [[#Constant2]]
17+
18+
; ModuleID = 'source.cpp'
19+
source_filename = "source.cpp"
20+
target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
21+
target triple = "spir"
22+
23+
@var = dso_local addrspace(2) constant i32 1, align 4
24+
@var2 = dso_local addrspace(2) constant i32 3, align 4
25+
@var.const = private unnamed_addr addrspace(2) constant i32 1, align 4
26+
@var2.const = private unnamed_addr addrspace(2) constant i32 3, align 4
27+
28+
; Function Attrs: convergent noinline norecurse nounwind optnone
29+
define dso_local spir_kernel void @test() #0 !kernel_arg_addr_space !2 !kernel_arg_access_qual !2 !kernel_arg_type !2 !kernel_arg_base_type !2 !kernel_arg_type_qual !2 !kernel_arg_host_accessible !2 !kernel_arg_pipe_depth !2 !kernel_arg_pipe_io !2 !kernel_arg_buffer_location !2 {
30+
entry:
31+
%0 = load i32, i32 addrspace(2)* @var.const, align 4
32+
%1 = load i32, i32 addrspace(2)* @var2.const, align 4
33+
%mul = mul nsw i32 %0, %1
34+
%mul1 = mul nsw i32 %mul, 2
35+
ret void
36+
}
37+
38+
attributes #0 = { convergent noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="false" }
39+
40+
!opencl.enable.FP_CONTRACT = !{}
41+
!opencl.ocl.version = !{!0}
42+
!opencl.spir.version = !{!0}
43+
!llvm.module.flags = !{!1}
44+
!opencl.used.extensions = !{!2}
45+
!opencl.used.optional.core.features = !{!2}
46+
!opencl.compiler.options = !{!2}
47+
!llvm.ident = !{!3}
48+
49+
!0 = !{i32 2, i32 0}
50+
!1 = !{i32 7, !"frame-pointer", i32 2}
51+
!2 = !{}
52+
!3 = !{!"Compiler"}

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_inline_assembly/inline_asm_clobbers.cl

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -20,7 +20,7 @@ size_t __ovld __cnfn get_global_id(unsigned int dimindx);
2020
// XCHECK-LLVM: [[STRUCTYPE:%[a-z0-9]+]] = type { i32, i32 }
2121

2222
// CHECK-LLVM-LABEL: define spir_kernel void @mem_clobber
23-
// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} """~{cc},~{memory}"
23+
// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "" "~{cc},~{memory}"
2424
// CHECK-LLVM: [[VALUE:%[0-9]+]] = load ptr addrspace(1), ptr
2525
// CHECK-LLVM-NEXT: getelementptr inbounds i32, ptr addrspace(1) [[VALUE]], i64 0
2626
// CHECK-LLVM-NEXT: store i32 1, ptr addrspace(1)
@@ -34,7 +34,7 @@ kernel void mem_clobber(global int *x) {
3434
}
3535

3636
// CHECK-LLVM-LABEL: define spir_kernel void @out_clobber
37-
// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "earlyclobber_instruction_out $0""=&r"
37+
// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "earlyclobber_instruction_out $0" "=&r"
3838
// CHECK-LLVM: barrier
3939
// CHECK-LLVM: store i32 %{{[a-z0-9]+}}, ptr [[VALUE:%[a-z0-9]+]], align 4
4040
// CHECK-LLVM-NEXT: [[STOREVAL:%[a-z0-9]+]] = call i32 asm "earlyclobber_instruction_out $0", "=&r"()
@@ -54,7 +54,7 @@ kernel void out_clobber(global int *x) {
5454
// Or bug in clang FE. To investigate later, change xchecks to checks and enable
5555

5656
// XCHECK-LLVM-LABEL: define spir_kernel void @in_clobber
57-
// XCHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "earlyclobber_instruction_in $0""&r"
57+
// XCHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "earlyclobber_instruction_in $0" "&r"
5858
// XCHECK-LLVM: barrier
5959
// XCHECK-LLVM: getelementptr
6060
// XCHECK-LLVM: store i32 %{{[a-z0-9]+}}, ptr [[LOADVAL:%[a-z0-9]+]], align 4
@@ -74,7 +74,7 @@ kernel void in_clobber(global int *x) {
7474
#endif
7575

7676
// XCHECK-LLVM-LABEL: define spir_kernel void @mixed_clobber
77-
// XCHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixedclobber_instruction $0 $1 $2""=&r,=&r,&r,1,~{cc},~{memory}"
77+
// XCHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixedclobber_instruction $0 $1 $2" "=&r,=&r,&r,1,~{cc},~{memory}"
7878

7979
#if 0
8080
kernel void mixed_clobber(global int *x, global int *y, global int *z) {

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_inline_assembly/inline_asm_constraints.cl

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -24,7 +24,7 @@ size_t __ovld __cnfn get_global_id(unsigned int dimindx);
2424
// CHECK-LLVM: [[STRUCTYPE:%[a-z]+]] = type { i32, i8, float }
2525

2626
// CHECK-LLVM-LABEL: define spir_kernel void @test_int
27-
// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "intcommand $0 $1""=r,r"
27+
// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "intcommand $0 $1" "=r,r"
2828
// CHECK-LLVM: [[VALUE:%[0-9]+]] = call i32 asm sideeffect "intcommand $0 $1", "=r,r"(i32 %{{[0-9]+}})
2929
// CHECK-LLVM-NEXT: store i32 [[VALUE]], ptr addrspace(1)
3030

@@ -34,7 +34,7 @@ kernel void test_int(global int *in, global int *out) {
3434
}
3535

3636
// CHECK-LLVM-LABEL: define spir_kernel void @test_float
37-
// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "floatcommand $0 $1""=r,r"
37+
// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "floatcommand $0 $1" "=r,r"
3838
// CHECK-LLVM: [[VALUE:%[0-9]+]] = call float asm sideeffect "floatcommand $0 $1", "=r,r"(float %{{[0-9]+}})
3939
// CHECK-LLVM-NEXT: store float [[VALUE]], ptr addrspace(1)
4040

@@ -44,7 +44,7 @@ kernel void test_float(global float *in, global float *out) {
4444
}
4545

4646
// CHECK-LLVM-LABEL: define spir_kernel void @test_mixed_integral
47-
// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixed_integral_command $0 $3 $1 $2""=r,r,r,r"
47+
// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixed_integral_command $0 $3 $1 $2" "=r,r,r,r"
4848
// CHECK-LLVM: [[VALUE:%[0-9]+]] = call i64 asm sideeffect "mixed_integral_command $0 $3 $1 $2", "=r,r,r,r"(i16 %{{[0-9]+}}, i32 %{{[0-9]+}}, i8 %{{[0-9]+}})
4949
// CHECK-LLVM-NEXT: store i64 [[VALUE]], ptr addrspace(1)
5050

@@ -55,7 +55,7 @@ kernel void test_mixed_integral(global uchar *A, global ushort *B, global uint *
5555
}
5656

5757
// CHECK-LLVM-LABEL: define spir_kernel void @test_mixed_floating
58-
// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixed_floating_command $0 $1 $2""=r,r,r"
58+
// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixed_floating_command $0 $1 $2" "=r,r,r"
5959
// CHECK-LLVM: [[VALUE:%[0-9]+]] = call half asm sideeffect "mixed_floating_command $0 $1 $2", "=r,r,r"(double %{{[0-9]+}}, float %{{[0-9]+}})
6060
// CHECK-LLVM-NEXT: store half [[VALUE]], ptr addrspace(1)
6161

@@ -66,7 +66,7 @@ kernel void test_mixed_floating(global float *A, global half *B, global double *
6666
}
6767

6868
// CHECK-LLVM-LABEL: define spir_kernel void @test_mixed_all
69-
// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixed_all_command $0 $3 $1 $2""=r,r,r,r"
69+
// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixed_all_command $0 $3 $1 $2" "=r,r,r,r"
7070
// CHECK-LLVM: [[VALUE:%[0-9]+]] = call i8 asm sideeffect "mixed_all_command $0 $3 $1 $2", "=r,r,r,r"(float %{{[0-9]+}}, i32 %{{[0-9]+}}, i8 %{{[0-9]+}})
7171
// CHECK-LLVM-NEXT: store i8 [[VALUE]], ptr addrspace(1)
7272

@@ -77,7 +77,7 @@ kernel void test_mixed_all(global uchar *A, global float *B, global uint *C, glo
7777
}
7878

7979
// CHECK-LLVM-LABEL: define spir_kernel void @test_multiple
80-
// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "multiple_command $0 $0 $1 $1 $2 $2""=r,=r,=r,0,1,2"
80+
// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "multiple_command $0 $0 $1 $1 $2 $2" "=r,=r,=r,0,1,2"
8181
// CHECK-LLVM: [[VALUE:%[0-9]+]] = call [[STRUCTYPE]] asm sideeffect "multiple_command $0 $0 $1 $1 $2 $2", "=r,=r,=r,0,1,2"(i32 %{{[0-9]+}}, i8 %{{[0-9]+}}, float %{{[0-9]+}})
8282
// CHECK-LLVM-NEXT: extractvalue [[STRUCTYPE]] [[VALUE]], 0
8383
// CHECK-LLVM-NEXT: extractvalue [[STRUCTYPE]] [[VALUE]], 1
@@ -90,7 +90,7 @@ kernel void test_multiple(global uchar *A, global float *B, global uint *C) {
9090
}
9191

9292
// CHECK-LLVM-LABEL: define spir_kernel void @test_constants
93-
// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "constcommand $0 $1""i,i"
93+
// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "constcommand $0 $1" "i,i"
9494
// CHECK-LLVM: call void asm sideeffect "constcommand $0 $1", "i,i"(i32 1, double 2.000000e+00)
9595

9696
kernel void test_constants() {

0 commit comments

Comments
 (0)