Lines Matching refs:Kern
631 msgpack::MapDocNode Kern) { in emitKernelLanguage() argument
640 Kern[".language"] = Kern.getDocument()->getNode("OpenCL C"); in emitKernelLanguage()
641 auto LanguageVersion = Kern.getDocument()->getArrayNode(); in emitKernelLanguage()
642 LanguageVersion.push_back(Kern.getDocument()->getNode( in emitKernelLanguage()
644 LanguageVersion.push_back(Kern.getDocument()->getNode( in emitKernelLanguage()
646 Kern[".language_version"] = LanguageVersion; in emitKernelLanguage()
650 msgpack::MapDocNode Kern) { in emitKernelAttrs() argument
653 Kern[".reqd_workgroup_size"] = getWorkGroupDimensions(Node); in emitKernelAttrs()
655 Kern[".workgroup_size_hint"] = getWorkGroupDimensions(Node); in emitKernelAttrs()
657 Kern[".vec_type_hint"] = Kern.getDocument()->getNode( in emitKernelAttrs()
664 Kern[".device_enqueue_symbol"] = Kern.getDocument()->getNode( in emitKernelAttrs()
671 msgpack::MapDocNode Kern) { in emitKernelArgs() argument
679 Kern[".args"] = Args; in emitKernelArgs()
856 auto Kern = HSAMetadataDoc->getMapNode(); in getHSAKernelProps() local
859 Kern[".kernarg_segment_size"] = Kern.getDocument()->getNode( in getHSAKernelProps()
861 Kern[".group_segment_fixed_size"] = in getHSAKernelProps()
862 Kern.getDocument()->getNode(ProgramInfo.LDSSize); in getHSAKernelProps()
863 Kern[".private_segment_fixed_size"] = in getHSAKernelProps()
864 Kern.getDocument()->getNode(ProgramInfo.ScratchSize); in getHSAKernelProps()
865 Kern[".kernarg_segment_align"] = in getHSAKernelProps()
866 Kern.getDocument()->getNode(std::max(Align(4), MaxKernArgAlign).value()); in getHSAKernelProps()
867 Kern[".wavefront_size"] = in getHSAKernelProps()
868 Kern.getDocument()->getNode(STM.getWavefrontSize()); in getHSAKernelProps()
869 Kern[".sgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.NumSGPR); in getHSAKernelProps()
870 Kern[".vgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.NumVGPR); in getHSAKernelProps()
871 Kern[".max_flat_workgroup_size"] = in getHSAKernelProps()
872 Kern.getDocument()->getNode(MFI.getMaxFlatWorkGroupSize()); in getHSAKernelProps()
873 Kern[".sgpr_spill_count"] = in getHSAKernelProps()
874 Kern.getDocument()->getNode(MFI.getNumSpilledSGPRs()); in getHSAKernelProps()
875 Kern[".vgpr_spill_count"] = in getHSAKernelProps()
876 Kern.getDocument()->getNode(MFI.getNumSpilledVGPRs()); in getHSAKernelProps()
878 return Kern; in getHSAKernelProps()
906 auto Kern = getHSAKernelProps(MF, ProgramInfo); in emitKernel() local
915 Kern[".name"] = Kern.getDocument()->getNode(Func.getName()); in emitKernel()
916 Kern[".symbol"] = Kern.getDocument()->getNode( in emitKernel()
918 emitKernelLanguage(Func, Kern); in emitKernel()
919 emitKernelAttrs(Func, Kern); in emitKernel()
920 emitKernelArgs(Func, Kern); in emitKernel()
923 Kernels.push_back(Kern); in emitKernel()