[llvm] [SPIR-V]: add SPIR-V extension: SPV_INTEL_variable_length_array (PR #83002)
Vyacheslav Levytskyy via llvm-commits
llvm-commits at lists.llvm.org
Mon Feb 26 05:43:29 PST 2024
https://github.com/VyacheslavLevytskyy created https://github.com/llvm/llvm-project/pull/83002
This PR adds SPIR-V extension SPV_INTEL_variable_length_array that allows to allocate local arrays whose number of elements is unknown at compile time:
* add a new SPIR-V internal intrinsic:int_spv_alloca_array
* legalize G_STACKSAVE and G_STACKRESTORE
* implement allocation of arrays (previously getArraySize() of AllocaInst was not used)
* add tests
>From bf54c255127ad9636255e730b3600ffca8bad42d Mon Sep 17 00:00:00 2001
From: "Levytskyy, Vyacheslav" <vyacheslav.levytskyy at intel.com>
Date: Mon, 26 Feb 2024 05:40:50 -0800
Subject: [PATCH] add SPIR-V extension: SPV_INTEL_variable_length_array
---
llvm/include/llvm/IR/IntrinsicsSPIRV.td | 1 +
llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp | 18 ++-
llvm/lib/Target/SPIRV/SPIRVInstrInfo.td | 9 ++
.../Target/SPIRV/SPIRVInstructionSelector.cpp | 56 +++++++++
llvm/lib/Target/SPIRV/SPIRVLegalizerInfo.cpp | 2 +
llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp | 8 ++
llvm/lib/Target/SPIRV/SPIRVSubtarget.cpp | 4 +
.../lib/Target/SPIRV/SPIRVSymbolicOperands.td | 2 +
.../SPV_INTEL_variable_length_array/vararr.ll | 54 +++++++++
.../vararr_spec_const.ll | 110 ++++++++++++++++++
10 files changed, 263 insertions(+), 1 deletion(-)
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_variable_length_array/vararr.ll
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_variable_length_array/vararr_spec_const.ll
diff --git a/llvm/include/llvm/IR/IntrinsicsSPIRV.td b/llvm/include/llvm/IR/IntrinsicsSPIRV.td
index 057dc64e88c26e..d6eabc5d24079a 100644
--- a/llvm/include/llvm/IR/IntrinsicsSPIRV.td
+++ b/llvm/include/llvm/IR/IntrinsicsSPIRV.td
@@ -33,6 +33,7 @@ let TargetPrefix = "spv" in {
def int_spv_cmpxchg : Intrinsic<[llvm_i32_ty], [llvm_any_ty, llvm_vararg_ty]>;
def int_spv_unreachable : Intrinsic<[], []>;
def int_spv_alloca : Intrinsic<[llvm_any_ty], []>;
+ def int_spv_alloca_array : Intrinsic<[llvm_any_ty], [llvm_anyint_ty]>;
def int_spv_undef : Intrinsic<[llvm_i32_ty], []>;
// Expect, Assume Intrinsics
diff --git a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
index e32cd50be56e38..afb24bfb322391 100644
--- a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
@@ -500,9 +500,25 @@ Instruction *SPIRVEmitIntrinsics::visitStoreInst(StoreInst &I) {
}
Instruction *SPIRVEmitIntrinsics::visitAllocaInst(AllocaInst &I) {
+ Value *ArraySize = nullptr;
+ if (I.isArrayAllocation()) {
+ const SPIRVSubtarget *STI = TM->getSubtargetImpl(*I.getFunction());
+ if (!STI->canUseExtension(
+ SPIRV::Extension::SPV_INTEL_variable_length_array))
+ report_fatal_error(
+ "array allocation: this instruction requires the following "
+ "SPIR-V extension: SPV_INTEL_variable_length_array",
+ false);
+ ArraySize = I.getArraySize();
+ }
+
TrackConstants = false;
Type *PtrTy = I.getType();
- auto *NewI = IRB->CreateIntrinsic(Intrinsic::spv_alloca, {PtrTy}, {});
+ auto *NewI =
+ ArraySize
+ ? IRB->CreateIntrinsic(Intrinsic::spv_alloca_array,
+ {PtrTy, ArraySize->getType()}, {ArraySize})
+ : IRB->CreateIntrinsic(Intrinsic::spv_alloca, {PtrTy}, {});
std::string InstName = I.hasName() ? I.getName().str() : "";
I.replaceAllUsesWith(NewI);
I.eraseFromParent();
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstrInfo.td b/llvm/lib/Target/SPIRV/SPIRVInstrInfo.td
index 7c5252e8cb372b..fe8c909236cde3 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstrInfo.td
+++ b/llvm/lib/Target/SPIRV/SPIRVInstrInfo.td
@@ -287,6 +287,15 @@ def OpPtrNotEqual: Op<402, (outs ID:$res), (ins TYPE:$resType, ID:$a, ID:$b),
def OpPtrDiff: Op<403, (outs ID:$res), (ins TYPE:$resType, ID:$a, ID:$b),
"$res = OpPtrDiff $resType $a $b">;
+// - SPV_INTEL_variable_length_array
+
+def OpVariableLengthArrayINTEL: Op<5818, (outs ID:$res), (ins TYPE:$type, ID:$length),
+ "$res = OpVariableLengthArrayINTEL $type $length">;
+def OpSaveMemoryINTEL: Op<5819, (outs ID:$res), (ins TYPE:$type),
+ "$res = OpSaveMemoryINTEL $type">;
+def OpRestoreMemoryINTEL: Op<5820, (outs), (ins ID:$ptr),
+ "OpRestoreMemoryINTEL $ptr">;
+
// 3.42.9 Function Instructions
def OpFunction: Op<54, (outs ID:$func),
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index 7258d3b4d88ed3..eb6dfbcc0cd780 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -99,6 +99,10 @@ class SPIRVInstructionSelector : public InstructionSelector {
MachineInstr &I) const;
bool selectStore(MachineInstr &I) const;
+ bool selectStackSave(Register ResVReg, const SPIRVType *ResType,
+ MachineInstr &I) const;
+ bool selectStackRestore(MachineInstr &I) const;
+
bool selectMemOperation(Register ResVReg, MachineInstr &I) const;
bool selectAtomicRMW(Register ResVReg, const SPIRVType *ResType,
@@ -165,6 +169,8 @@ class SPIRVInstructionSelector : public InstructionSelector {
bool selectFrameIndex(Register ResVReg, const SPIRVType *ResType,
MachineInstr &I) const;
+ bool selectAllocaArray(Register ResVReg, const SPIRVType *ResType,
+ MachineInstr &I) const;
bool selectBranch(MachineInstr &I) const;
bool selectBranchCond(MachineInstr &I) const;
@@ -504,6 +510,11 @@ bool SPIRVInstructionSelector::spvSelect(Register ResVReg,
case TargetOpcode::G_FENCE:
return selectFence(I);
+ case TargetOpcode::G_STACKSAVE:
+ return selectStackSave(ResVReg, ResType, I);
+ case TargetOpcode::G_STACKRESTORE:
+ return selectStackRestore(I);
+
default:
return false;
}
@@ -649,6 +660,35 @@ bool SPIRVInstructionSelector::selectStore(MachineInstr &I) const {
return MIB.constrainAllUses(TII, TRI, RBI);
}
+bool SPIRVInstructionSelector::selectStackSave(Register ResVReg,
+ const SPIRVType *ResType,
+ MachineInstr &I) const {
+ if (!STI.canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
+ report_fatal_error(
+ "llvm.stacksave intrinsic: this instruction requires the following "
+ "SPIR-V extension: SPV_INTEL_variable_length_array",
+ false);
+ MachineBasicBlock &BB = *I.getParent();
+ return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSaveMemoryINTEL))
+ .addDef(ResVReg)
+ .addUse(GR.getSPIRVTypeID(ResType))
+ .constrainAllUses(TII, TRI, RBI);
+}
+
+bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &I) const {
+ if (!STI.canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
+ report_fatal_error(
+ "llvm.stackrestore intrinsic: this instruction requires the following "
+ "SPIR-V extension: SPV_INTEL_variable_length_array",
+ false);
+ if (!I.getOperand(0).isReg())
+ return false;
+ MachineBasicBlock &BB = *I.getParent();
+ return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpRestoreMemoryINTEL))
+ .addUse(I.getOperand(0).getReg())
+ .constrainAllUses(TII, TRI, RBI);
+}
+
bool SPIRVInstructionSelector::selectMemOperation(Register ResVReg,
MachineInstr &I) const {
MachineBasicBlock &BB = *I.getParent();
@@ -1461,6 +1501,8 @@ bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
break;
case Intrinsic::spv_alloca:
return selectFrameIndex(ResVReg, ResType, I);
+ case Intrinsic::spv_alloca_array:
+ return selectAllocaArray(ResVReg, ResType, I);
case Intrinsic::spv_assume:
if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_expect_assume))
BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpAssumeTrueKHR))
@@ -1480,6 +1522,20 @@ bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
return true;
}
+bool SPIRVInstructionSelector::selectAllocaArray(Register ResVReg,
+ const SPIRVType *ResType,
+ MachineInstr &I) const {
+ // there was an allocation size parameter to the allocation instruction
+ // that is not 1
+ MachineBasicBlock &BB = *I.getParent();
+ return BuildMI(BB, I, I.getDebugLoc(),
+ TII.get(SPIRV::OpVariableLengthArrayINTEL))
+ .addDef(ResVReg)
+ .addUse(GR.getSPIRVTypeID(ResType))
+ .addUse(I.getOperand(2).getReg())
+ .constrainAllUses(TII, TRI, RBI);
+}
+
bool SPIRVInstructionSelector::selectFrameIndex(Register ResVReg,
const SPIRVType *ResType,
MachineInstr &I) const {
diff --git a/llvm/lib/Target/SPIRV/SPIRVLegalizerInfo.cpp b/llvm/lib/Target/SPIRV/SPIRVLegalizerInfo.cpp
index 4f2e7a240fc2cc..1937cb64c206e3 100644
--- a/llvm/lib/Target/SPIRV/SPIRVLegalizerInfo.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVLegalizerInfo.cpp
@@ -186,6 +186,8 @@ SPIRVLegalizerInfo::SPIRVLegalizerInfo(const SPIRVSubtarget &ST) {
getActionDefinitionsBuilder(G_IMPLICIT_DEF).alwaysLegal();
+ getActionDefinitionsBuilder({G_STACKSAVE, G_STACKRESTORE}).alwaysLegal();
+
getActionDefinitionsBuilder(G_INTTOPTR)
.legalForCartesianProduct(allPtrs, allIntScalars);
getActionDefinitionsBuilder(G_PTRTOINT)
diff --git a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
index 3be28c97d95381..ac3d6b362d350b 100644
--- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
@@ -1110,6 +1110,14 @@ void addInstrRequirements(const MachineInstr &MI,
case SPIRV::OpAtomicFMaxEXT:
AddAtomicFloatRequirements(MI, Reqs, ST);
break;
+ case SPIRV::OpVariableLengthArrayINTEL:
+ case SPIRV::OpSaveMemoryINTEL:
+ case SPIRV::OpRestoreMemoryINTEL:
+ if (ST.canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array)) {
+ Reqs.addExtension(SPIRV::Extension::SPV_INTEL_variable_length_array);
+ Reqs.addCapability(SPIRV::Capability::VariableLengthArrayINTEL);
+ }
+ break;
default:
break;
}
diff --git a/llvm/lib/Target/SPIRV/SPIRVSubtarget.cpp b/llvm/lib/Target/SPIRV/SPIRVSubtarget.cpp
index 79f16146ccd944..0e8952dc6a9c9f 100644
--- a/llvm/lib/Target/SPIRV/SPIRVSubtarget.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVSubtarget.cpp
@@ -85,6 +85,10 @@ cl::list<SPIRV::Extension::Extension> Extensions(
"SPV_KHR_subgroup_rotate",
"Adds a new instruction that enables rotating values across "
"invocations within a subgroup."),
+ clEnumValN(SPIRV::Extension::SPV_INTEL_variable_length_array,
+ "SPV_INTEL_variable_length_array",
+ "Allows to allocate local arrays whose number of elements "
+ "is unknown at compile time."),
clEnumValN(SPIRV::Extension::SPV_INTEL_function_pointers,
"SPV_INTEL_function_pointers",
"Allows translation of function pointers.")));
diff --git a/llvm/lib/Target/SPIRV/SPIRVSymbolicOperands.td b/llvm/lib/Target/SPIRV/SPIRVSymbolicOperands.td
index b022b97408d7d4..211c22340eb82c 100644
--- a/llvm/lib/Target/SPIRV/SPIRVSymbolicOperands.td
+++ b/llvm/lib/Target/SPIRV/SPIRVSymbolicOperands.td
@@ -296,6 +296,7 @@ defm SPV_INTEL_fpga_latency_control : ExtensionOperand<101>;
defm SPV_INTEL_fpga_argument_interfaces : ExtensionOperand<102>;
defm SPV_INTEL_optnone : ExtensionOperand<103>;
defm SPV_INTEL_function_pointers : ExtensionOperand<104>;
+defm SPV_INTEL_variable_length_array : ExtensionOperand<105>;
//===----------------------------------------------------------------------===//
// Multiclass used to define Capabilities enum values and at the same time
@@ -462,6 +463,7 @@ defm AtomicFloat16AddEXT : CapabilityOperand<6095, 0, 0, [SPV_EXT_shader_atomic_
defm AtomicFloat16MinMaxEXT : CapabilityOperand<5616, 0, 0, [SPV_EXT_shader_atomic_float_min_max], []>;
defm AtomicFloat32MinMaxEXT : CapabilityOperand<5612, 0, 0, [SPV_EXT_shader_atomic_float_min_max], []>;
defm AtomicFloat64MinMaxEXT : CapabilityOperand<5613, 0, 0, [SPV_EXT_shader_atomic_float_min_max], []>;
+defm VariableLengthArrayINTEL : CapabilityOperand<5817, 0, 0, [SPV_INTEL_variable_length_array], []>;
defm GroupUniformArithmeticKHR : CapabilityOperand<6400, 0, 0, [SPV_KHR_uniform_group_instructions], []>;
defm USMStorageClassesINTEL : CapabilityOperand<5935, 0, 0, [SPV_INTEL_usm_storage_classes], [Kernel]>;
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_variable_length_array/vararr.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_variable_length_array/vararr.ll
new file mode 100644
index 00000000000000..897aab70852d2e
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_variable_length_array/vararr.ll
@@ -0,0 +1,54 @@
+; Modified from: https://github.com/KhronosGroup/SPIRV-LLVM-Translator/test/extensions/INTEL/SPV_INTEL_variable_length_array/basic.ll
+
+; RUN: not llc -O0 -mtriple=spirv32-unknown-unknown %s -o %t.spvt 2>&1 | FileCheck %s --check-prefix=CHECK-ERROR
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-extensions=SPV_INTEL_variable_length_array %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown --spirv-extensions=SPV_INTEL_variable_length_array %s -o - -filetype=obj | spirv-val %}
+
+; CHECK-ERROR: LLVM ERROR: array allocation: this instruction requires the following SPIR-V extension: SPV_INTEL_variable_length_array
+
+; CHECK-SPIRV: Capability VariableLengthArrayINTEL
+; CHECK-SPIRV: Extension "SPV_INTEL_variable_length_array"
+
+; CHECK-SPIRV-DAG: OpName %[[Len:.*]] "a"
+; CHECK-SPIRV-DAG: %[[Long:.*]] = OpTypeInt 64 0
+; CHECK-SPIRV-DAG: %[[Int:.*]] = OpTypeInt 32 0
+; CHECK-SPIRV-DAG: %[[Char:.*]] = OpTypeInt 8 0
+; CHECK-SPIRV-DAG: %[[CharPtr:.*]] = OpTypePointer {{[a-zA-Z]+}} %[[Char]]
+; CHECK-SPIRV-DAG: %[[IntPtr:.*]] = OpTypePointer {{[a-zA-Z]+}} %[[Int]]
+; CHECK-SPIRV: %[[Len]] = OpFunctionParameter %[[Long:.*]]
+; CHECK-SPIRV: %[[SavedMem1:.*]] = OpSaveMemoryINTEL %[[CharPtr]]
+; CHECK-SPIRV: OpVariableLengthArrayINTEL %[[IntPtr]] %[[Len]]
+; CHECK-SPIRV: OpRestoreMemoryINTEL %[[SavedMem1]]
+; CHECK-SPIRV: %[[SavedMem2:.*]] = OpSaveMemoryINTEL %[[CharPtr]]
+; CHECK-SPIRV: OpVariableLengthArrayINTEL %[[IntPtr]] %[[Len]]
+; CHECK-SPIRV: OpRestoreMemoryINTEL %[[SavedMem2]]
+
+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"
+target triple = "spir"
+
+define dso_local spir_func i32 @foo(i64 %a, i64 %b) {
+entry:
+ %vector1 = alloca [42 x i32], align 16
+ call void @llvm.lifetime.start.p0(i64 168, ptr nonnull %vector1)
+ %stack1 = call ptr @llvm.stacksave.p0()
+ %vla = alloca i32, i64 %a, align 16
+ %arrayidx = getelementptr inbounds i32, ptr %vla, i64 %b
+ %elem1 = load i32, ptr %arrayidx, align 4
+ call void @llvm.stackrestore.p0(ptr %stack1)
+ %stack2 = call ptr @llvm.stacksave.p0()
+ %vla2 = alloca i32, i64 %a, align 16
+ %arrayidx3 = getelementptr inbounds [42 x i32], ptr %vector1, i64 0, i64 %b
+ %elemt = load i32, ptr %arrayidx3, align 4
+ %add = add nsw i32 %elemt, %elem1
+ %arrayidx4 = getelementptr inbounds i32, ptr %vla2, i64 %b
+ %elem2 = load i32, ptr %arrayidx4, align 4
+ %add5 = add nsw i32 %add, %elem2
+ call void @llvm.stackrestore.p0(ptr %stack2)
+ call void @llvm.lifetime.end.p0(i64 168, ptr nonnull %vector1)
+ ret i32 %add5
+}
+
+declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture)
+declare ptr @llvm.stacksave.p0()
+declare void @llvm.stackrestore.p0(ptr)
+declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture)
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_variable_length_array/vararr_spec_const.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_variable_length_array/vararr_spec_const.ll
new file mode 100644
index 00000000000000..bf090aed2a0ac2
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_variable_length_array/vararr_spec_const.ll
@@ -0,0 +1,110 @@
+; Modified from: https://github.com/KhronosGroup/SPIRV-LLVM-Translator/test/extensions/INTEL/SPV_INTEL_variable_length_array/vla_spec_const.ll
+
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-extensions=SPV_INTEL_variable_length_array %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown --spirv-extensions=SPV_INTEL_variable_length_array %s -o - -filetype=obj | spirv-val %}
+
+; CHECK-SPIRV: Capability VariableLengthArrayINTEL
+; CHECK-SPIRV: Extension "SPV_INTEL_variable_length_array"
+; CHECK-SPIRV: OpDecorate %[[SpecConst:.*]] SpecId 0
+; CHECK-SPIRV-DAG: %[[Long:.*]] = OpTypeInt 64 0
+; CHECK-SPIRV-DAG: %[[Int:.*]] = OpTypeInt 32 0
+; CHECK-SPIRV-DAG: %[[IntPtr:.*]] = OpTypePointer {{[a-zA-Z]+}} %[[Int]]
+; CHECK-SPIRV: %[[SpecConst]] = OpSpecConstant %[[Long]]
+; CHECK-SPIRV-LABEL: FunctionEnd
+; CHECK-SPIRV: %[[SpecConstVal:.*]] = OpFunctionCall %[[Long]]
+; CHECK-SPIRV: OpSaveMemoryINTEL
+; CHECK-SPIRV: OpVariableLengthArrayINTEL %[[IntPtr]] %[[SpecConstVal]]
+; CHECK-SPIRV: OpRestoreMemoryINTEL
+
+; CHECK-SPIRV: OpFunction %[[Long]]
+; CHECK-SPIRV: ReturnValue %[[SpecConst]]
+
+target datalayout = "e-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 = "spir64-unknown-linux"
+
+%"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon" = type { %"class._ZTSN2cl4sycl12experimental13spec_constantIm13MyUInt64ConstEE.cl::sycl::experimental::spec_constant" }
+%"class._ZTSN2cl4sycl12experimental13spec_constantIm13MyUInt64ConstEE.cl::sycl::experimental::spec_constant" = type { i8 }
+
+$_ZTS17SpecializedKernel = comdat any
+
+$_ZNK2cl4sycl12experimental13spec_constantIm13MyUInt64ConstE3getEv = comdat any
+
+; Function Attrs: norecurse
+define weak_odr dso_local spir_kernel void @_ZTS17SpecializedKernel() #0 comdat !kernel_arg_addr_space !4 !kernel_arg_access_qual !4 !kernel_arg_type !4 !kernel_arg_base_type !4 !kernel_arg_type_qual !4 {
+entry:
+ %0 = alloca %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon", align 1
+ call void @llvm.lifetime.start.p0(i64 1, ptr %0) #4
+ %1 = addrspacecast ptr %0 to ptr addrspace(4)
+ call spir_func void @"_ZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlvE_clEv"(ptr addrspace(4) %1)
+ call void @llvm.lifetime.end.p0(i64 1, ptr %0) #4
+ ret void
+}
+
+; Function Attrs: argmemonly nounwind willreturn
+declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1
+
+; Function Attrs: inlinehint norecurse
+define internal spir_func void @"_ZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlvE_clEv"(ptr addrspace(4) %this) #2 align 2 {
+entry:
+ %this.addr = alloca ptr addrspace(4), align 8
+ %saved_stack = alloca ptr, align 8
+ %__vla_expr0 = alloca i64, align 8
+ store ptr addrspace(4) %this, ptr %this.addr, align 8, !tbaa !5
+ %this1 = load ptr addrspace(4), ptr %this.addr, align 8
+ %call = call spir_func i64 @_ZNK2cl4sycl12experimental13spec_constantIm13MyUInt64ConstE3getEv(ptr addrspace(4) %this1)
+ %0 = call ptr @llvm.stacksave.p0()
+ store ptr %0, ptr %saved_stack, align 8
+ %vla = alloca i32, i64 %call, align 4
+ store i64 %call, ptr %__vla_expr0, align 8
+ store i32 42, ptr %vla, align 4, !tbaa !9
+ %1 = load ptr, ptr %saved_stack, align 8
+ call void @llvm.stackrestore.p0(ptr %1)
+ ret void
+}
+
+; Function Attrs: argmemonly nounwind willreturn
+declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1
+
+; Function Attrs: norecurse
+define linkonce_odr dso_local spir_func i64 @_ZNK2cl4sycl12experimental13spec_constantIm13MyUInt64ConstE3getEv(ptr addrspace(4) %this) #3 comdat align 2 {
+entry:
+ %this.addr = alloca ptr addrspace(4), align 8
+ %TName = alloca ptr addrspace(4), align 8
+ store ptr addrspace(4) %this, ptr %this.addr, align 8, !tbaa !5
+ call void @llvm.lifetime.start.p0(i64 8, ptr %TName) #4
+ %0 = call i64 @_Z20__spirv_SpecConstantix(i32 0, i64 0), !SYCL_SPEC_CONST_SYM_ID !11
+ call void @llvm.lifetime.end.p0(i64 8, ptr %TName) #4
+ ret i64 %0
+}
+
+; Function Attrs: nounwind
+declare ptr @llvm.stacksave.p0() #4
+
+; Function Attrs: nounwind
+declare void @llvm.stackrestore.p0(ptr) #4
+
+declare i64 @_Z20__spirv_SpecConstantix(i32, i64)
+
+attributes #0 = { norecurse "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="/work/intel/vla_spec_const.cpp" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
+attributes #1 = { argmemonly nounwind willreturn }
+attributes #2 = { inlinehint norecurse "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
+attributes #3 = { norecurse "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
+attributes #4 = { nounwind }
+
+!llvm.module.flags = !{!0}
+!opencl.spir.version = !{!1}
+!spirv.Source = !{!2}
+!llvm.ident = !{!3}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 1, i32 2}
+!2 = !{i32 4, i32 100000}
+!3 = !{!"clang version 12.0.0"}
+!4 = !{}
+!5 = !{!6, !6, i64 0}
+!6 = !{!"any pointer", !7, i64 0}
+!7 = !{!"omnipotent char", !8, i64 0}
+!8 = !{!"Simple C++ TBAA"}
+!9 = !{!10, !10, i64 0}
+!10 = !{!"int", !7, i64 0}
+!11 = !{!"_ZTS13MyUInt64Const", i32 0}
More information about the llvm-commits
mailing list