diff options
Diffstat (limited to 'llvm')
10 files changed, 263 insertions, 1 deletions
diff --git a/llvm/include/llvm/IR/IntrinsicsSPIRV.td b/llvm/include/llvm/IR/IntrinsicsSPIRV.td index 057dc64..d6eabc5 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 e32cd50..afb24bf 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 7c5252e..fe8c909 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 f1e18f0..9b38073 100644 --- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp +++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp @@ -99,6 +99,10 @@ private: 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, @@ -167,6 +171,8 @@ private: 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; @@ -508,6 +514,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; } @@ -653,6 +664,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(); @@ -1505,6 +1545,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)) @@ -1524,6 +1566,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 aedca79..049ca4a 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, G_FREEZE}).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 3be28c9..ac3d6b3 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 79f1614..0e8952d 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 b022b97..211c223 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 0000000..897aab7 --- /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 0000000..fbac43e --- /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: + %p = alloca %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon", align 1 + call void @llvm.lifetime.start.p0(i64 1, ptr %p) #4 + %p4 = addrspacecast ptr %p to ptr addrspace(4) + call spir_func void @"_ZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlvE_clEv"(ptr addrspace(4) %p4) + call void @llvm.lifetime.end.p0(i64 1, ptr %p) #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) + %p = call ptr @llvm.stacksave.p0() + store ptr %p, 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 + %torestore = load ptr, ptr %saved_stack, align 8 + call void @llvm.stackrestore.p0(ptr %torestore) + 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 + %p = 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 %p +} + +; 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} |