From 64db1408abc92a32b8da757a2b68b5d58a422b57 Mon Sep 17 00:00:00 2001 From: Diana Chen Date: Wed, 27 Sep 2023 00:37:20 +0000 Subject: [PATCH] IGC: Fixed LSC checker in VectorProcess In VectorProcess::reLayoutLoadStore, checks if LSC is used by "EmitPass::shouldGenerateLSCQuery" instead of "CPlatform::LSCEnabled" to align with the checker used in EmitPass::emitStoreRawIndexed. This avoids QW scatter/gather messages being generated on platforms those supporting LSC but don't support QW scatter/gather messages. --- IGC/Compiler/CISACodeGen/VectorProcess.cpp | 2 +- .../tests/VectorProcess/bindless_store.ll | 447 ++++++++++++++++++ 2 files changed, 448 insertions(+), 1 deletion(-) create mode 100644 IGC/Compiler/tests/VectorProcess/bindless_store.ll diff --git a/IGC/Compiler/CISACodeGen/VectorProcess.cpp b/IGC/Compiler/CISACodeGen/VectorProcess.cpp index db8271d010ef..91770d439e29 100644 --- a/IGC/Compiler/CISACodeGen/VectorProcess.cpp +++ b/IGC/Compiler/CISACodeGen/VectorProcess.cpp @@ -311,7 +311,7 @@ bool VectorProcess::reLayoutLoadStore(Instruction* Inst) useQW = has_QW_BTS_GS && nelts == 1 && (eTyBytes == 8U && align >= 8U); } - if (cgCtx->platform.LSCEnabled()) + if (EmitPass::shouldGenerateLSCQuery(*cgCtx, Inst) == Tristate::True) { // With LSC, want to use QW if element size is 8 bytes. useQW = (eTyBytes == 8); diff --git a/IGC/Compiler/tests/VectorProcess/bindless_store.ll b/IGC/Compiler/tests/VectorProcess/bindless_store.ll new file mode 100644 index 000000000000..681c5b2a1b58 --- /dev/null +++ b/IGC/Compiler/tests/VectorProcess/bindless_store.ll @@ -0,0 +1,447 @@ +;=========================== begin_copyright_notice ============================ +; +; Copyright (C) 2023 Intel Corporation +; +; SPDX-License-Identifier: MIT +; +;============================ end_copyright_notice ============================= +; This test checks that i64 data write is translated to v2i32 in VectorProcess. +; Ref: IGC-7889 + +; Below .ll is generated by the ocloc options +; -device mtl +; -internal_options "-cl-intel-use-bindless-mode -cl-intel-use-bindless-advanced-mode" +; CL Source: +; __kernel void QueryKernelTimestamps(__global ulong* dst) { +; uint gid = get_global_id(0); +; uint currentOffset = gid * 4; +; dst[currentOffset] = 0; +; } +; +; RUN: igc_opt -enable-debugify --igc-vectorprocess -platformmtl -S < %s 2>&1 | FileCheck %s + +; ------------------------------------------------ +; OCL_asmfbcbaf4428005986_0186_CG_after_Removeredundantinstructions.ll +; ------------------------------------------------ +target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024-n8:16:32" +target triple = "spir64-unknown-unknown" + +; Function Attrs: convergent nounwind +define spir_kernel void @QueryKernelTimestamps(i64 addrspace(1)* %dst, <8 x i32> %r0, <8 x i32> %payloadHeader, <3 x i32> %enqueuedLocalSize, i16 %localIdX, i16 %localIdY, i16 %localIdZ, i32 %bufferOffset) #0 { +entry: + %payloadHeader.scalar = extractelement <8 x i32> %payloadHeader, i32 0 + %enqueuedLocalSize.scalar = extractelement <3 x i32> %enqueuedLocalSize, i32 0 + %r0.scalar17 = extractelement <8 x i32> %r0, i32 1 + %mul.i.i.i = mul i32 %enqueuedLocalSize.scalar, %r0.scalar17 + %localIdX2 = zext i16 %localIdX to i32 + %add.i.i.i = add i32 %mul.i.i.i, %localIdX2 + %add4.i.i.i = add i32 %add.i.i.i, %payloadHeader.scalar + %mul = shl i32 %add4.i.i.i, 5 + %0 = addrspacecast i64 addrspace(1)* %dst to i64 addrspace(2490368)* + ; CHECK: call void @llvm.genx.GenISA.storerawvector + ; CHECK-NOT: i64 + ; CHECK-SAME: v2i32.v2i32 + call void @llvm.genx.GenISA.storeraw.indexed.p2490368i64.i64(i64 addrspace(2490368)* %0, i32 %mul, i64 0, i32 8, i1 false) + ret void +} + +; Function Attrs: convergent nounwind readnone +declare spir_func i32 @__builtin_IB_get_group_id(i32) local_unnamed_addr #1 + +; Function Attrs: convergent nounwind readnone +declare spir_func i32 @__builtin_IB_get_enqueued_local_size(i32) local_unnamed_addr #1 + +; Function Attrs: convergent nounwind readnone +declare spir_func i32 @__builtin_IB_get_local_id_x() local_unnamed_addr #1 + +; Function Attrs: nounwind +declare void @llvm.assume(i1) #2 + +; Function Attrs: convergent nounwind readnone +declare spir_func i32 @__builtin_IB_get_global_offset(i32) local_unnamed_addr #1 + +; Function Attrs: convergent nounwind readnone +declare spir_func i32 @__builtin_IB_get_local_id_y() local_unnamed_addr #1 + +; Function Attrs: convergent nounwind readnone +declare spir_func i32 @__builtin_IB_get_local_id_z() local_unnamed_addr #1 + +; Function Attrs: argmemonly nounwind writeonly +declare void @llvm.genx.GenISA.storeraw.indexed.p2490368i64.i64(i64 addrspace(2490368)*, i32, i64, i32, i1) #3 + +attributes #0 = { convergent nounwind "less-precise-fpmad"="true" } +attributes #1 = { convergent nounwind readnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #2 = { nounwind } +attributes #3 = { argmemonly nounwind writeonly } + +!IGCMetadata = !{!0} +!igc.functions = !{!349} +!opencl.ocl.version = !{!361, !361, !361, !361, !361} +!opencl.spir.version = !{!361, !361, !361, !361, !361} +!llvm.ident = !{!362, !362, !362, !362, !362} +!llvm.module.flags = !{!363} + +!0 = !{!"ModuleMD", !1, !2, !83, !170, !201, !202, !206, !207, !208, !240, !262, !275, !276, !277, !290, !291, !292, !293, !294, !295, !296, !297, !298, !299, !303, !304, !305, !306, !307, !308, !309, !310, !311, !312, !313, !314, !315, !316, !317, !319, !323, !324, !325, !326, !327, !328, !329, !330, !331, !332, !333, !334, !335, !336, !337, !338, !339, !146, !340, !341, !342, !344, !347, !348} +!1 = !{!"isPrecise", i1 false} +!2 = !{!"compOpt", !3, !4, !5, !6, !7, !8, !9, !10, !11, !12, !13, !14, !15, !16, !17, !18, !19, !20, !21, !22, !23, !24, !25, !26, !27, !28, !29, !30, !31, !32, !33, !34, !35, !36, !37, !38, !39, !40, !41, !42, !43, !44, !45, !46, !47, !48, !49, !50, !51, !52, !53, !54, !55, !56, !57, !58, !59, !60, !61, !62, !63, !64, !65, !66, !67, !68, !69, !70, !71, !72, !73, !74, !75, !76, !77, !78, !79, !80, !81, !82} +!3 = !{!"DenormsAreZero", i1 false} +!4 = !{!"BFTFDenormsAreZero", i1 false} +!5 = !{!"CorrectlyRoundedDivSqrt", i1 false} +!6 = !{!"OptDisable", i1 false} +!7 = !{!"MadEnable", i1 true} +!8 = !{!"NoSignedZeros", i1 false} +!9 = !{!"NoNaNs", i1 false} +!10 = !{!"FloatRoundingMode", i32 0} +!11 = !{!"FloatCvtIntRoundingMode", i32 3} +!12 = !{!"LoadCacheDefault", i32 4} +!13 = !{!"StoreCacheDefault", i32 2} +!14 = !{!"VISAPreSchedRPThreshold", i32 0} +!15 = !{!"SetLoopUnrollThreshold", i32 0} +!16 = !{!"UnsafeMathOptimizations", i1 false} +!17 = !{!"disableCustomUnsafeOpts", i1 false} +!18 = !{!"disableReducePow", i1 false} +!19 = !{!"FiniteMathOnly", i1 false} +!20 = !{!"FastRelaxedMath", i1 false} +!21 = !{!"DashGSpecified", i1 false} +!22 = !{!"FastCompilation", i1 false} +!23 = !{!"UseScratchSpacePrivateMemory", i1 true} +!24 = !{!"RelaxedBuiltins", i1 false} +!25 = !{!"SubgroupIndependentForwardProgressRequired", i1 true} +!26 = !{!"GreaterThan2GBBufferRequired", i1 true} +!27 = !{!"GreaterThan4GBBufferRequired", i1 false} +!28 = !{!"DisableA64WA", i1 false} +!29 = !{!"ForceEnableA64WA", i1 false} +!30 = !{!"PushConstantsEnable", i1 true} +!31 = !{!"HasPositivePointerOffset", i1 false} +!32 = !{!"HasBufferOffsetArg", i1 true} +!33 = !{!"BufferOffsetArgOptional", i1 true} +!34 = !{!"HasSubDWAlignedPtrArg", i1 false} +!35 = !{!"replaceGlobalOffsetsByZero", i1 false} +!36 = !{!"forcePixelShaderSIMDMode", i32 0} +!37 = !{!"pixelShaderDoNotAbortOnSpill", i1 false} +!38 = !{!"UniformWGS", i1 false} +!39 = !{!"disableVertexComponentPacking", i1 false} +!40 = !{!"disablePartialVertexComponentPacking", i1 false} +!41 = !{!"PreferBindlessImages", i1 true} +!42 = !{!"UseBindlessMode", i1 true} +!43 = !{!"UseLegacyBindlessMode", i1 true} +!44 = !{!"disableMathRefactoring", i1 false} +!45 = !{!"atomicBranch", i1 false} +!46 = !{!"ForceInt32DivRemEmu", i1 false} +!47 = !{!"ForceInt32DivRemEmuSP", i1 false} +!48 = !{!"WaveIntrinsicUsed", i1 false} +!49 = !{!"DisableMultiPolyPS", i1 false} +!50 = !{!"NeedTexture3DLODWA", i1 false} +!51 = !{!"DisableFastestSingleCSSIMD", i1 false} +!52 = !{!"DisableFastestLinearScan", i1 false} +!53 = !{!"UseStatelessforPrivateMemory", i1 false} +!54 = !{!"EnableTakeGlobalAddress", i1 false} +!55 = !{!"IsLibraryCompilation", i1 false} +!56 = !{!"FastVISACompile", i1 false} +!57 = !{!"MatchSinCosPi", i1 false} +!58 = !{!"ExcludeIRFromZEBinary", i1 false} +!59 = !{!"EmitZeBinVISASections", i1 false} +!60 = !{!"FP64GenEmulationEnabled", i1 false} +!61 = !{!"allowDisableRematforCS", i1 false} +!62 = !{!"DisableIncSpillCostAllAddrTaken", i1 false} +!63 = !{!"DisableCPSOmaskWA", i1 false} +!64 = !{!"DisableFastestGopt", i1 false} +!65 = !{!"WaForceHalfPromotionComputeShader", i1 false} +!66 = !{!"WaForceHalfPromotionPixelVertexShader", i1 false} +!67 = !{!"DisableConstantCoalescing", i1 false} +!68 = !{!"EnableUndefAlphaOutputAsRed", i1 true} +!69 = !{!"WaEnableALTModeVisaWA", i1 false} +!70 = !{!"WaEnableAtomicWaveFusion", i1 false} +!71 = !{!"WaEnableAtomicWaveFusionNonNullResource", i1 false} +!72 = !{!"WaEnableAtomicWaveFusionStateless", i1 false} +!73 = !{!"WaEnableAtomicWaveFusionTyped", i1 false} +!74 = !{!"NewSpillCostFunction", i1 false} +!75 = !{!"EnableVRT", i1 false} +!76 = !{!"ForceLargeGRFNum4RQ", i1 false} +!77 = !{!"EnableURBWritesMerging", i1 true} +!78 = !{!"DisableEUFusion", i1 false} +!79 = !{!"DisableFDivToFMulInvOpt", i1 false} +!80 = !{!"initializePhiSampleSourceWA", i1 false} +!81 = !{!"WaDisableSubspanUseNoMaskForCB", i1 false} +!82 = !{!"FastestS1Options", i32 0} +!83 = !{!"FuncMD", !84, !85} +!84 = !{!"FuncMDMap[0]", void (i64 addrspace(1)*, <8 x i32>, <8 x i32>, <3 x i32>, i16, i16, i16, i32)* @QueryKernelTimestamps} +!85 = !{!"FuncMDValue[0]", !86, !87, !91, !92, !93, !94, !95, !96, !97, !119, !138, !139, !140, !141, !142, !143, !144, !145, !146, !147, !148, !149, !150, !151, !152, !153, !155, !157, !159, !161, !163, !165, !166} +!86 = !{!"localOffsets"} +!87 = !{!"workGroupWalkOrder", !88, !89, !90} +!88 = !{!"dim0", i32 0} +!89 = !{!"dim1", i32 0} +!90 = !{!"dim2", i32 0} +!91 = !{!"funcArgs"} +!92 = !{!"functionType", !"KernelFunction"} +!93 = !{!"inlineDynConstants"} +!94 = !{!"inlineDynRootConstant"} +!95 = !{!"inlineDynConstantDescTable"} +!96 = !{!"m_pInterestingConstants"} +!97 = !{!"rtInfo", !98, !99, !100, !101, !102, !103, !104, !105, !106, !107, !108, !109, !110, !111, !112, !113, !117, !118, !75} +!98 = !{!"callableShaderType", !"NumberOfCallableShaderTypes"} +!99 = !{!"isContinuation", i1 false} +!100 = !{!"hasTraceRayPayload", i1 false} +!101 = !{!"hasHitAttributes", i1 false} +!102 = !{!"hasCallableData", i1 false} +!103 = !{!"ShaderStackSize", i32 0} +!104 = !{!"ShaderHash", i64 0} +!105 = !{!"ShaderName", !""} +!106 = !{!"ParentName", !""} +!107 = !{!"SlotNum", i1* null} +!108 = !{!"NOSSize", i32 0} +!109 = !{!"globalRootSignatureSize", i32 0} +!110 = !{!"Entries"} +!111 = !{!"SpillUnions"} +!112 = !{!"CustomHitAttrSizeInBytes", i32 0} +!113 = !{!"Types", !114, !115, !116} +!114 = !{!"FrameStartTys"} +!115 = !{!"ArgumentTys"} +!116 = !{!"FullFrameTys"} +!117 = !{!"Aliases"} +!118 = !{!"NumGRF", i32 0} +!119 = !{!"resAllocMD", !120, !121, !122, !123, !137} +!120 = !{!"uavsNumType", i32 1} +!121 = !{!"srvsNumType", i32 0} +!122 = !{!"samplersNumType", i32 0} +!123 = !{!"argAllocMDList", !124, !128, !131, !132, !133, !134, !135, !136} +!124 = !{!"argAllocMDListVec[0]", !125, !126, !127} +!125 = !{!"type", i32 1} +!126 = !{!"extensionType", i32 -1} +!127 = !{!"indexType", i32 0} +!128 = !{!"argAllocMDListVec[1]", !129, !126, !130} +!129 = !{!"type", i32 0} +!130 = !{!"indexType", i32 -1} +!131 = !{!"argAllocMDListVec[2]", !129, !126, !130} +!132 = !{!"argAllocMDListVec[3]", !129, !126, !130} +!133 = !{!"argAllocMDListVec[4]", !129, !126, !130} +!134 = !{!"argAllocMDListVec[5]", !129, !126, !130} +!135 = !{!"argAllocMDListVec[6]", !129, !126, !130} +!136 = !{!"argAllocMDListVec[7]", !129, !126, !130} +!137 = !{!"inlineSamplersMD"} +!138 = !{!"maxByteOffsets"} +!139 = !{!"IsInitializer", i1 false} +!140 = !{!"IsFinalizer", i1 false} +!141 = !{!"CompiledSubGroupsNumber", i32 0} +!142 = !{!"hasInlineVmeSamplers", i1 false} +!143 = !{!"localSize", i32 0} +!144 = !{!"localIDPresent", i1 false} +!145 = !{!"groupIDPresent", i1 false} +!146 = !{!"privateMemoryPerWI", i32 0} +!147 = !{!"globalIDPresent", i1 false} +!148 = !{!"hasSyncRTCalls", i1 false} +!149 = !{!"hasNonKernelArgLoad", i1 false} +!150 = !{!"hasNonKernelArgStore", i1 false} +!151 = !{!"hasNonKernelArgAtomic", i1 false} +!152 = !{!"UserAnnotations"} +!153 = !{!"m_OpenCLArgAddressSpaces", !154} +!154 = !{!"m_OpenCLArgAddressSpacesVec[0]", i32 1} +!155 = !{!"m_OpenCLArgAccessQualifiers", !156} +!156 = !{!"m_OpenCLArgAccessQualifiersVec[0]", !"none"} +!157 = !{!"m_OpenCLArgTypes", !158} +!158 = !{!"m_OpenCLArgTypesVec[0]", !"ulong*"} +!159 = !{!"m_OpenCLArgBaseTypes", !160} +!160 = !{!"m_OpenCLArgBaseTypesVec[0]", !"long*"} +!161 = !{!"m_OpenCLArgTypeQualifiers", !162} +!162 = !{!"m_OpenCLArgTypeQualifiersVec[0]", !""} +!163 = !{!"m_OpenCLArgNames", !164} +!164 = !{!"m_OpenCLArgNamesVec[0]", !"dst"} +!165 = !{!"m_OpenCLArgScalarAsPointers"} +!166 = !{!"m_OptsToDisablePerFunc", !167, !168, !169} +!167 = !{!"m_OptsToDisablePerFuncSet[0]", !"IGC-AddressArithmeticSinking"} +!168 = !{!"m_OptsToDisablePerFuncSet[1]", !"IGC-AllowSimd32Slicing"} +!169 = !{!"m_OptsToDisablePerFuncSet[2]", !"IGC-SinkLoadOpt"} +!170 = !{!"pushInfo", !171, !172, !173, !177, !178, !179, !180, !181, !182, !183, !184, !197, !198, !199, !200} +!171 = !{!"pushableAddresses"} +!172 = !{!"bindlessPushInfo"} +!173 = !{!"dynamicBufferInfo", !174, !175, !176} +!174 = !{!"firstIndex", i32 0} +!175 = !{!"numOffsets", i32 0} +!176 = !{!"forceDisabled", i1 false} +!177 = !{!"MaxNumberOfPushedBuffers", i32 0} +!178 = !{!"inlineConstantBufferSlot", i32 -1} +!179 = !{!"inlineConstantBufferOffset", i32 -1} +!180 = !{!"inlineConstantBufferGRFOffset", i32 -1} +!181 = !{!"constants"} +!182 = !{!"inputs"} +!183 = !{!"constantReg"} +!184 = !{!"simplePushInfoArr", !185, !194, !195, !196} +!185 = !{!"simplePushInfoArrVec[0]", !186, !187, !188, !189, !190, !191, !192, !193} +!186 = !{!"cbIdx", i32 0} +!187 = !{!"pushableAddressGrfOffset", i32 -1} +!188 = !{!"pushableOffsetGrfOffset", i32 -1} +!189 = !{!"offset", i32 0} +!190 = !{!"size", i32 0} +!191 = !{!"isStateless", i1 false} +!192 = !{!"isBindless", i1 false} +!193 = !{!"simplePushLoads"} +!194 = !{!"simplePushInfoArrVec[1]", !186, !187, !188, !189, !190, !191, !192, !193} +!195 = !{!"simplePushInfoArrVec[2]", !186, !187, !188, !189, !190, !191, !192, !193} +!196 = !{!"simplePushInfoArrVec[3]", !186, !187, !188, !189, !190, !191, !192, !193} +!197 = !{!"simplePushBufferUsed", i32 0} +!198 = !{!"pushAnalysisWIInfos"} +!199 = !{!"inlineRTGlobalPtrOffset", i32 0} +!200 = !{!"rtSyncSurfPtrOffset", i32 0} +!201 = !{!"WaEnableICBPromotion", i1 false} +!202 = !{!"vsInfo", !203, !204, !205} +!203 = !{!"DrawIndirectBufferIndex", i32 -1} +!204 = !{!"vertexReordering", i32 -1} +!205 = !{!"MaxNumOfOutputs", i32 0} +!206 = !{!"dsInfo", !205} +!207 = !{!"gsInfo", !205} +!208 = !{!"psInfo", !209, !210, !211, !212, !213, !214, !215, !216, !217, !218, !219, !220, !221, !222, !223, !224, !225, !226, !227, !228, !229, !230, !231, !232, !233, !234, !235, !236, !237, !238, !239} +!209 = !{!"BlendStateDisabledMask", i8 0} +!210 = !{!"SkipSrc0Alpha", i1 false} +!211 = !{!"DualSourceBlendingDisabled", i1 false} +!212 = !{!"ForceEnableSimd32", i1 false} +!213 = !{!"outputDepth", i1 false} +!214 = !{!"outputStencil", i1 false} +!215 = !{!"outputMask", i1 false} +!216 = !{!"blendToFillEnabled", i1 false} +!217 = !{!"forceEarlyZ", i1 false} +!218 = !{!"hasVersionedLoop", i1 false} +!219 = !{!"forceSingleSourceRTWAfterDualSourceRTW", i1 false} +!220 = !{!"requestCPSizeRelevant", i1 false} +!221 = !{!"requestCPSize", i1 false} +!222 = !{!"texelMaskFastClearMode", !"Disabled"} +!223 = !{!"NumSamples", i8 0} +!224 = !{!"blendOptimizationMode"} +!225 = !{!"colorOutputMask"} +!226 = !{!"ProvokingVertexModeNosIndex", i32 0} +!227 = !{!"ProvokingVertexModeNosPatch", !""} +!228 = !{!"ProvokingVertexModeLast", !"Negative"} +!229 = !{!"VertexAttributesBypass", i1 false} +!230 = !{!"LegacyBaryAssignmentDisableLinear", i1 false} +!231 = !{!"LegacyBaryAssignmentDisableLinearNoPerspective", i1 false} +!232 = !{!"LegacyBaryAssignmentDisableLinearCentroid", i1 false} +!233 = !{!"LegacyBaryAssignmentDisableLinearNoPerspectiveCentroid", i1 false} +!234 = !{!"LegacyBaryAssignmentDisableLinearSample", i1 false} +!235 = !{!"LegacyBaryAssignmentDisableLinearNoPerspectiveSample", i1 false} +!236 = !{!"MeshShaderWAPerPrimitiveUserDataEnable", !"Negative"} +!237 = !{!"meshShaderWAPerPrimitiveUserDataEnablePatchName", !""} +!238 = !{!"generatePatchesForRTWriteSends", i1 false} +!239 = !{!"WaDisableVRS", i1 false} +!240 = !{!"csInfo", !241, !242, !243, !244, !245, !14, !15, !246, !247, !248, !249, !250, !251, !252, !253, !254, !255, !256, !45, !257, !258, !259, !260, !261} +!241 = !{!"maxWorkGroupSize", i32 0} +!242 = !{!"waveSize", i32 0} +!243 = !{!"ComputeShaderSecondCompile"} +!244 = !{!"forcedSIMDSize", i8 0} +!245 = !{!"forceTotalGRFNum", i32 0} +!246 = !{!"allowLowerSimd", i1 false} +!247 = !{!"disableSimd32Slicing", i1 false} +!248 = !{!"disableSplitOnSpill", i1 false} +!249 = !{!"enableNewSpillCostFunction", i1 false} +!250 = !{!"forcedVISAPreRAScheduler", i1 false} +!251 = !{!"forceUniformBuffer", i1 false} +!252 = !{!"forceUniformSurfaceSampler", i1 false} +!253 = !{!"disableLocalIdOrderOptimizations", i1 false} +!254 = !{!"disableDispatchAlongY", i1 false} +!255 = !{!"neededThreadIdLayout", i1* null} +!256 = !{!"forceTileYWalk", i1 false} +!257 = !{!"walkOrderEnabled", i1 false} +!258 = !{!"walkOrderOverride", i32 0} +!259 = !{!"ResForHfPacking"} +!260 = !{!"hasWaveMatrix", i1 false} +!261 = !{!"constantFoldSimdSize", i1 false} +!262 = !{!"msInfo", !263, !264, !265, !266, !267, !268, !269, !270, !271, !272, !273, !228, !226, !274} +!263 = !{!"PrimitiveTopology", i32 3} +!264 = !{!"MaxNumOfPrimitives", i32 0} +!265 = !{!"MaxNumOfVertices", i32 0} +!266 = !{!"MaxNumOfPerPrimitiveOutputs", i32 0} +!267 = !{!"MaxNumOfPerVertexOutputs", i32 0} +!268 = !{!"WorkGroupSize", i32 0} +!269 = !{!"WorkGroupMemorySizeInBytes", i32 0} +!270 = !{!"IndexFormat", i32 6} +!271 = !{!"SubgroupSize", i32 0} +!272 = !{!"VPandRTAIndexAutostripEnable", i1 false} +!273 = !{!"MeshShaderWAPerPrimitiveUserDataEnable", i1 false} +!274 = !{!"numPrimitiveAttributesPatchBaseName", !""} +!275 = !{!"taskInfo", !205, !268, !269, !271} +!276 = !{!"NBarrierCnt", i32 0} +!277 = !{!"rtInfo", !278, !279, !280, !281, !282, !283, !284, !285, !286, !287, !288, !289} +!278 = !{!"RayQueryAllocSizeInBytes", i32 0} +!279 = !{!"NumContinuations", i32 0} +!280 = !{!"RTAsyncStackAddrspace", i32 -1} +!281 = !{!"RTAsyncStackSurfaceStateOffset", i1* null} +!282 = !{!"SWHotZoneAddrspace", i32 -1} +!283 = !{!"SWHotZoneSurfaceStateOffset", i1* null} +!284 = !{!"SWStackAddrspace", i32 -1} +!285 = !{!"SWStackSurfaceStateOffset", i1* null} +!286 = !{!"RTSyncStackAddrspace", i32 -1} +!287 = !{!"RTSyncStackSurfaceStateOffset", i1* null} +!288 = !{!"doSyncDispatchRays", i1 false} +!289 = !{!"MemStyle", !"Xe"} +!290 = !{!"EnableTextureIndirection", i1 false} +!291 = !{!"EnableSamplerIndirection", i1 false} +!292 = !{!"samplerStateStride", i32 0} +!293 = !{!"samplerStateOffset", i32 0} +!294 = !{!"textureStateStride", i32 0} +!295 = !{!"textureStateOffset", i32 0} +!296 = !{!"CurUniqueIndirectIdx", i32 0} +!297 = !{!"inlineDynTextures"} +!298 = !{!"inlineResInfoData"} +!299 = !{!"immConstant", !300, !301, !302} +!300 = !{!"data"} +!301 = !{!"sizes"} +!302 = !{!"zeroIdxs"} +!303 = !{!"stringConstants"} +!304 = !{!"inlineConstantBuffers"} +!305 = !{!"inlineGlobalBuffers"} +!306 = !{!"GlobalPointerProgramBinaryInfos"} +!307 = !{!"ConstantPointerProgramBinaryInfos"} +!308 = !{!"GlobalBufferAddressRelocInfo"} +!309 = !{!"ConstantBufferAddressRelocInfo"} +!310 = !{!"forceLscCacheList"} +!311 = !{!"SrvMap"} +!312 = !{!"RootConstantBufferOffsetInBytes"} +!313 = !{!"RasterizerOrderedByteAddressBuffer"} +!314 = !{!"RasterizerOrderedViews"} +!315 = !{!"MinNOSPushConstantSize", i32 0} +!316 = !{!"inlineProgramScopeOffsets"} +!317 = !{!"shaderData", !318} +!318 = !{!"numReplicas", i32 0} +!319 = !{!"URBInfo", !320, !321, !322} +!320 = !{!"has64BVertexHeaderInput", i1 false} +!321 = !{!"has64BVertexHeaderOutput", i1 false} +!322 = !{!"hasVertexHeader", i1 true} +!323 = !{!"UseBindlessImage", i1 true} +!324 = !{!"enableRangeReduce", i1 false} +!325 = !{!"disableNewTrigFuncRangeReduction", i1 false} +!326 = !{!"enableFRemToSRemOpt", i1 false} +!327 = !{!"enableSampleptrToLdmsptrSample0", i1 false} +!328 = !{!"WaForceSIMD32MicropolyRasterize", i1 false} +!329 = !{!"WaEnableFastestForAllWaveIntrinsicsCS", i1 false} +!330 = !{!"WaEnableFastestForAllWaveIntrinsicsPS", i1 false} +!331 = !{!"allowMatchMadOptimizationforVS", i1 false} +!332 = !{!"disableMatchMadOptimizationForCS", i1 false} +!333 = !{!"disableMemOptforNegativeOffsetLoads", i1 false} +!334 = !{!"enableThreeWayLoadSpiltOpt", i1 false} +!335 = !{!"statefulResourcesNotAliased", i1 false} +!336 = !{!"disableMixMode", i1 false} +!337 = !{!"genericAccessesResolved", i1 false} +!338 = !{!"disableSeparateSpillPvtScratchSpace", i1 false} +!339 = !{!"disableSeparateScratchWA", i1 false} +!340 = !{!"PrivateMemoryPerFG"} +!341 = !{!"m_OptsToDisable"} +!342 = !{!"capabilities", !343} +!343 = !{!"globalVariableDecorationsINTEL", i1 false} +!344 = !{!"m_ShaderResourceViewMcsMask", !345, !346} +!345 = !{!"m_ShaderResourceViewMcsMaskVec[0]", i64 0} +!346 = !{!"m_ShaderResourceViewMcsMaskVec[1]", i64 0} +!347 = !{!"computedDepthMode", i32 0} +!348 = !{!"isHDCFastClearShader", i1 false} +!349 = !{void (i64 addrspace(1)*, <8 x i32>, <8 x i32>, <3 x i32>, i16, i16, i16, i32)* @QueryKernelTimestamps, !350} +!350 = !{!351, !352} +!351 = !{!"function_type", i32 0} +!352 = !{!"implicit_arg_desc", !353, !354, !355, !356, !357, !358, !359} +!353 = !{i32 0} +!354 = !{i32 1} +!355 = !{i32 6} +!356 = !{i32 7} +!357 = !{i32 8} +!358 = !{i32 9} +!359 = !{i32 14, !360} +!360 = !{!"explicit_arg_num", i32 0} +!361 = !{i32 2, i32 0} +!362 = !{!"clang version 9.0.0 (e83f00113f5cbb67ea42dc66cdf1099e7c25e776)"} +!363 = !{i32 1, !"wchar_size", i32 4}