Skip to content

[llvm-spirv] Cherry pick Khronos changes to expand collection of entry point interfaces (PR #1334) #10623

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 2 commits into from
Aug 2, 2023
Merged

[llvm-spirv] Cherry pick Khronos changes to expand collection of entry point interfaces (PR #1334) #10623

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
d483857
Cherry pick 1334
asudarsa Jul 29, 2023
80a3686
Test changes from open-source
asudarsa Jul 30, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
30 changes: 18 additions & 12 deletions llvm-spirv/lib/SPIRV/SPIRVWriter.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -874,9 +874,7 @@ SPIRVFunction *LLVMToSPIRVBase::transFunctionDecl(Function *F) {

BM->setName(BF, F->getName().str());
}
if (isKernel(F))
BM->addEntryPoint(ExecutionModelKernel, BF->getId());
else if (F->getLinkage() != GlobalValue::InternalLinkage)
if (!isKernel(F) && F->getLinkage() != GlobalValue::InternalLinkage)
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We no longer add entry point for function decls.

BF->setLinkageType(transLinkageType(F));

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

void LLVMToSPIRVBase::collectInputOutputVariables(SPIRVFunction *SF,
Function *F) {
std::vector<SPIRVId>
LLVMToSPIRVBase::collectEntryPointInterfaces(SPIRVFunction *SF, Function *F) {
std::vector<SPIRVId> Interface;
for (auto &GV : M->globals()) {
const auto AS = GV.getAddressSpace();
if (AS != SPIRAS_Input && AS != SPIRAS_Output)
continue;
SPIRVModule *BM = SF->getModule();
if (!BM->isAllowedToUseVersion(VersionNumber::SPIRV_1_4))
if (AS != SPIRAS_Input && AS != SPIRAS_Output)
continue;

std::unordered_set<const Function *> Funcs;

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

if (isAnyFunctionReachableFromFunction(F, Funcs)) {
SF->addVariable(ValueMap[&GV]);
SPIRVWord ModuleVersion = static_cast<SPIRVWord>(BM->getSPIRVVersion());
if (AS != SPIRAS_Input && AS != SPIRAS_Output &&
ModuleVersion < static_cast<SPIRVWord>(VersionNumber::SPIRV_1_4))
BM->setMinSPIRVVersion(VersionNumber::SPIRV_1_4);
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We now pass Version number directly without converting to SPIRVWord

Interface.push_back(ValueMap[&GV]->getId());
}
}
return Interface;
}

void LLVMToSPIRVBase::mutateFuncArgType(
Expand Down Expand Up @@ -5118,10 +5124,10 @@ void LLVMToSPIRVBase::transFunction(Function *I) {
joinFPContract(I, FPContract::ENABLED);
fpContractUpdateRecursive(I, getFPContract(I));

bool IsKernelEntryPoint = isKernel(I);

if (IsKernelEntryPoint) {
collectInputOutputVariables(BF, I);
if (isKernel(I)) {
auto Interface = collectEntryPointInterfaces(BF, I);
BM->addEntryPoint(ExecutionModelKernel, BF->getId(), I->getName().str(),
Interface);
}
}

Expand Down
2 changes: 2 additions & 0 deletions llvm-spirv/lib/SPIRV/SPIRVWriter.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -248,6 +248,8 @@ class LLVMToSPIRVBase : protected BuiltinCallHelper {
const Function *FS,
const std::unordered_set<const Function *> Funcs) const;
void collectInputOutputVariables(SPIRVFunction *SF, Function *F);
std::vector<SPIRVId> collectEntryPointInterfaces(SPIRVFunction *BF,
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This function declaration was missing in the original patch.

Function *F);
};

class LLVMToSPIRVPass : public PassInfoMixin<LLVMToSPIRVPass> {
Expand Down
1 change: 0 additions & 1 deletion llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -268,7 +268,6 @@ class SPIRVDecorateLinkageAttr : public SPIRVDecorate {
#ifdef _SPIRV_SUPPORT_TEXT_FMT
if (SPIRVUseTextFormat) {
Encoder << getString(Literals.cbegin(), Literals.cend() - 1);
Encoder.OS << " ";
Encoder << (SPIRVLinkageTypeKind)Literals.back();
} else
#endif
Expand Down
6 changes: 4 additions & 2 deletions llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -621,9 +621,11 @@ void SPIRVEntryPoint::encode(spv_ostream &O) const {
}

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

void SPIRVExecutionMode::encode(spv_ostream &O) const {
Expand Down
1 change: 1 addition & 0 deletions llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -533,6 +533,7 @@ template <Op OC> class SPIRVAnnotation : public SPIRVAnnotationGeneric {

class SPIRVEntryPoint : public SPIRVAnnotation<OpEntryPoint> {
public:
static const SPIRVWord FixedWC = 4;
SPIRVEntryPoint(SPIRVModule *TheModule, SPIRVExecutionModelKind,
SPIRVId TheId, const std::string &TheName,
std::vector<SPIRVId> Variables);
Expand Down
38 changes: 12 additions & 26 deletions llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -128,20 +128,6 @@ class SPIRVModuleImpl : public SPIRVModule {
getValueTypes(const std::vector<SPIRVId> &) const override;
SPIRVMemoryModelKind getMemoryModel() const override { return MemoryModel; }
SPIRVConstant *getLiteralAsConstant(unsigned Literal) override;
unsigned getNumEntryPoints(SPIRVExecutionModelKind EM) const override {
auto Loc = EntryPointVec.find(EM);
if (Loc == EntryPointVec.end())
return 0;
return Loc->second.size();
}
SPIRVFunction *getEntryPoint(SPIRVExecutionModelKind EM,
unsigned I) const override {
auto Loc = EntryPointVec.find(EM);
if (Loc == EntryPointVec.end())
return nullptr;
assert(I < Loc->second.size());
return get<SPIRVFunction>(Loc->second[I]);
}
unsigned getNumFunctions() const override { return FuncVec.size(); }
unsigned getNumVariables() const override { return VariableVec.size(); }
SourceLanguage getSourceLanguage(SPIRVWord *Ver = nullptr) const override {
Expand Down Expand Up @@ -225,8 +211,9 @@ class SPIRVModuleImpl : public SPIRVModule {
SPIRVGroupMemberDecorate *
addGroupMemberDecorate(SPIRVDecorationGroup *Group,
const std::vector<SPIRVEntry *> &Targets) override;
void addEntryPoint(SPIRVExecutionModelKind ExecModel,
SPIRVId EntryPoint) override;
void addEntryPoint(SPIRVExecutionModelKind ExecModel, SPIRVId EntryPoint,
const std::string &Name,
const std::vector<SPIRVId> &Variables) override;
SPIRVForward *addForward(SPIRVType *Ty) override;
SPIRVForward *addForward(SPIRVId, SPIRVType *Ty) override;
SPIRVFunction *addFunction(SPIRVFunction *) override;
Expand Down Expand Up @@ -508,11 +495,11 @@ class SPIRVModuleImpl : public SPIRVModule {
typedef std::vector<SPIRVGroupDecorateGeneric *> SPIRVGroupDecVec;
typedef std::vector<SPIRVAsmTargetINTEL *> SPIRVAsmTargetVector;
typedef std::vector<SPIRVAsmINTEL *> SPIRVAsmVector;
typedef std::vector<SPIRVEntryPoint *> SPIRVEntryPointVec;
typedef std::map<SPIRVId, SPIRVExtInstSetKind> SPIRVIdToInstructionSetMap;
std::map<SPIRVExtInstSetKind, SPIRVId> ExtInstSetIds;
typedef std::map<SPIRVId, SPIRVExtInstSetKind> SPIRVIdToBuiltinSetMap;
typedef std::map<SPIRVExecutionModelKind, SPIRVIdSet> SPIRVExecModelIdSetMap;
typedef std::map<SPIRVExecutionModelKind, SPIRVIdVec> SPIRVExecModelIdVecMap;
typedef std::unordered_map<std::string, SPIRVString *> SPIRVStringMap;
typedef std::map<SPIRVTypeStruct *, std::vector<std::pair<unsigned, SPIRVId>>>
SPIRVUnknownStructFieldMap;
Expand Down Expand Up @@ -540,7 +527,7 @@ class SPIRVModuleImpl : public SPIRVModule {
SPIRVAsmTargetVector AsmTargetVec;
SPIRVAsmVector AsmVec;
SPIRVExecModelIdSetMap EntryPointSet;
SPIRVExecModelIdVecMap EntryPointVec;
SPIRVEntryPointVec EntryPointVec;
SPIRVStringMap StrMap;
SPIRVCapMap CapMap;
SPIRVUnknownStructFieldMap UnknownStructFieldMap;
Expand Down Expand Up @@ -1086,11 +1073,14 @@ SPIRVModuleImpl::addDecorate(SPIRVDecorateGeneric *Dec) {
}

void SPIRVModuleImpl::addEntryPoint(SPIRVExecutionModelKind ExecModel,
SPIRVId EntryPoint) {
SPIRVId EntryPoint, const std::string &Name,
const std::vector<SPIRVId> &Variables) {
assert(isValid(ExecModel) && "Invalid execution model");
assert(EntryPoint != SPIRVID_INVALID && "Invalid entry point");
auto *EP =
add(new SPIRVEntryPoint(this, ExecModel, EntryPoint, Name, Variables));
EntryPointVec.push_back(EP);
EntryPointSet[ExecModel].insert(EntryPoint);
EntryPointVec[ExecModel].push_back(EntryPoint);
addCapabilities(SPIRV::getCapability(ExecModel));
}

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

O << SPIRVMemoryModel(&M);

for (auto &I : MI.EntryPointVec)
for (auto &II : I.second)
O << SPIRVEntryPoint(&M, I.first, II, M.get<SPIRVFunction>(II)->getName(),
M.get<SPIRVFunction>(II)->getVariables());
O << MI.EntryPointVec;

for (auto &I : MI.EntryPointVec)
for (auto &II : I.second)
MI.get<SPIRVFunction>(II)->encodeExecutionModes(O);
MI.get<SPIRVFunction>(I->getTargetId())->encodeExecutionModes(O);

O << MI.StringVec;

Expand Down
7 changes: 3 additions & 4 deletions llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -133,14 +133,11 @@ class SPIRVModule {
virtual const SPIRVCapMap &getCapability() const = 0;
virtual bool hasCapability(SPIRVCapabilityKind) const = 0;
virtual SPIRVExtInstSetKind getBuiltinSet(SPIRVId) const = 0;
virtual SPIRVFunction *getEntryPoint(SPIRVExecutionModelKind,
unsigned) const = 0;
virtual std::set<std::string> &getExtension() = 0;
virtual SPIRVFunction *getFunction(unsigned) const = 0;
virtual SPIRVVariable *getVariable(unsigned) const = 0;
virtual SPIRVMemoryModelKind getMemoryModel() const = 0;
virtual unsigned getNumFunctions() const = 0;
virtual unsigned getNumEntryPoints(SPIRVExecutionModelKind) const = 0;
virtual unsigned getNumVariables() const = 0;
virtual SourceLanguage getSourceLanguage(SPIRVWord *) const = 0;
virtual std::set<std::string> &getSourceExtension() = 0;
Expand Down Expand Up @@ -223,7 +220,9 @@ class SPIRVModule {
const std::vector<SPIRVEntry *> &Targets) = 0;
virtual SPIRVGroupDecorateGeneric *
addGroupDecorateGeneric(SPIRVGroupDecorateGeneric *GDec) = 0;
virtual void addEntryPoint(SPIRVExecutionModelKind, SPIRVId) = 0;
virtual void addEntryPoint(SPIRVExecutionModelKind, SPIRVId,
const std::string &,
const std::vector<SPIRVId> &) = 0;
virtual SPIRVForward *addForward(SPIRVType *Ty) = 0;
virtual SPIRVForward *addForward(SPIRVId, SPIRVType *Ty) = 0;
virtual SPIRVFunction *addFunction(SPIRVFunction *) = 0;
Expand Down
1 change: 1 addition & 0 deletions llvm-spirv/lib/SPIRV/libSPIRV/SPIRVStream.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -180,6 +180,7 @@ const SPIRVEncoder &operator<<(const SPIRVEncoder &O, const std::string &Str) {
#ifdef _SPIRV_SUPPORT_TEXT_FMT
if (SPIRVUseTextFormat) {
writeQuotedString(O.OS, Str);
O.OS << " ";
return O;
}
#endif
Expand Down
3 changes: 0 additions & 3 deletions llvm-spirv/test/ExecutionMode.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,9 +1,6 @@
; RUN: llvm-as < %s | llvm-spirv -spirv-text -o %t
; RUN: FileCheck < %t %s

; check for magic number followed by version 1.1
; CHECK: 119734787 65792

; CHECK-DAG: TypeVoid [[VOID:[0-9]+]]

; CHECK-DAG: EntryPoint 6 [[WORKER:[0-9]+]] "worker"
Expand Down
2 changes: 1 addition & 1 deletion llvm-spirv/test/copy_object.spt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,7 +5,7 @@
2 Capability Int64
2 Capability Int8
3 MemoryModel 2 2
8 EntryPoint 6 1 "copy_object"
6 EntryPoint 6 1 "copy_object"
3 Source 3 102000
3 Name 2 "in"
4 Decorate 3 BuiltIn 28
Expand Down
52 changes: 52 additions & 0 deletions llvm-spirv/test/entry-point-interfaces.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,52 @@
; RUN: llvm-as %s -o %t.bc

; RUN: llvm-spirv %t.bc -o %t.spv
; RUN: spirv-val --target-env spv1.4 %t.spv
; RUN: llvm-spirv -to-text %t.spv -o %t.from.spv.spt
; RUN: FileCheck < %t.from.spv.spt %s --check-prefix=CHECK-SPIRV

; RUN: llvm-spirv -spirv-text %t.bc -o %t.from.bc.spt
; RUN: FileCheck < %t.from.bc.spt %s --check-prefix=CHECK-SPIRV

; CHECK-SPIRV: 7 EntryPoint 6 [[#]] "test" [[#Interface1:]] [[#Interface2:]]
; CHECK-SPIRV: TypeInt [[#TypeInt:]] 32 0
; CHECK-SPIRV: Constant [[#TypeInt]] [[#Constant1:]] 1
; CHECK-SPIRV: Constant [[#TypeInt]] [[#Constant2:]] 3
; CHECK-SPIRV: Variable [[#]] [[#Interface1]] 0 [[#Constant1]]
; CHECK-SPIRV: Variable [[#]] [[#Interface2]] 0 [[#Constant2]]

; ModuleID = 'source.cpp'
source_filename = "source.cpp"
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"
target triple = "spir"

@var = dso_local addrspace(2) constant i32 1, align 4
@var2 = dso_local addrspace(2) constant i32 3, align 4
@var.const = private unnamed_addr addrspace(2) constant i32 1, align 4
@var2.const = private unnamed_addr addrspace(2) constant i32 3, align 4

; Function Attrs: convergent noinline norecurse nounwind optnone
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 {
entry:
%0 = load i32, i32 addrspace(2)* @var.const, align 4
%1 = load i32, i32 addrspace(2)* @var2.const, align 4
%mul = mul nsw i32 %0, %1
%mul1 = mul nsw i32 %mul, 2
ret void
}

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" }

!opencl.enable.FP_CONTRACT = !{}
!opencl.ocl.version = !{!0}
!opencl.spir.version = !{!0}
!llvm.module.flags = !{!1}
!opencl.used.extensions = !{!2}
!opencl.used.optional.core.features = !{!2}
!opencl.compiler.options = !{!2}
!llvm.ident = !{!3}

!0 = !{i32 2, i32 0}
!1 = !{i32 7, !"frame-pointer", i32 2}
!2 = !{}
!3 = !{!"Compiler"}
8 changes: 4 additions & 4 deletions llvm-spirv/test/extensions/INTEL/SPV_INTEL_inline_assembly/inline_asm_clobbers.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,7 +20,7 @@ size_t __ovld __cnfn get_global_id(unsigned int dimindx);
// XCHECK-LLVM: [[STRUCTYPE:%[a-z0-9]+]] = type { i32, i32 }

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

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

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

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

#if 0
kernel void mixed_clobber(global int *x, global int *y, global int *z) {
Expand Down
14 changes: 7 additions & 7 deletions llvm-spirv/test/extensions/INTEL/SPV_INTEL_inline_assembly/inline_asm_constraints.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,7 +24,7 @@ size_t __ovld __cnfn get_global_id(unsigned int dimindx);
// CHECK-LLVM: [[STRUCTYPE:%[a-z]+]] = type { i32, i8, float }

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

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

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

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

// CHECK-LLVM-LABEL: define spir_kernel void @test_mixed_integral
// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixed_integral_command $0 $3 $1 $2""=r,r,r,r"
// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixed_integral_command $0 $3 $1 $2" "=r,r,r,r"
// 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]+}})
// CHECK-LLVM-NEXT: store i64 [[VALUE]], ptr addrspace(1)

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

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

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

// CHECK-LLVM-LABEL: define spir_kernel void @test_mixed_all
// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixed_all_command $0 $3 $1 $2""=r,r,r,r"
// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixed_all_command $0 $3 $1 $2" "=r,r,r,r"
// 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]+}})
// CHECK-LLVM-NEXT: store i8 [[VALUE]], ptr addrspace(1)

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

// CHECK-LLVM-LABEL: define spir_kernel void @test_multiple
// 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"
// 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"
// 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]+}})
// CHECK-LLVM-NEXT: extractvalue [[STRUCTYPE]] [[VALUE]], 0
// CHECK-LLVM-NEXT: extractvalue [[STRUCTYPE]] [[VALUE]], 1
Expand All @@ -90,7 +90,7 @@ kernel void test_multiple(global uchar *A, global float *B, global uint *C) {
}

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

kernel void test_constants() {
Expand Down
Loading