[SPIR-V]: add SPIR-V extension: SPV_INTEL_variable_length_array #83002
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-backend-spir-v @llvm/pr-subscribers-llvm-ir Author: Vyacheslav Levytskyy (VyacheslavLevytskyy) ChangesThis 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:
Patch is 20.51 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/83002.diff 10 Files Affected:
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"...
[truncated]
|
michalpaszkowski
approved these changes
Feb 27, 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 adds SPIR-V extension SPV_INTEL_variable_length_array that allows to allocate local arrays whose number of elements is unknown at compile time: