Lines Matching refs:Kern
215 msgpack::MapDocNode Kern) { in emitKernelLanguage() argument
224 Kern[".language"] = Kern.getDocument()->getNode("OpenCL C"); in emitKernelLanguage()
225 auto LanguageVersion = Kern.getDocument()->getArrayNode(); in emitKernelLanguage()
226 LanguageVersion.push_back(Kern.getDocument()->getNode( in emitKernelLanguage()
228 LanguageVersion.push_back(Kern.getDocument()->getNode( in emitKernelLanguage()
230 Kern[".language_version"] = LanguageVersion; in emitKernelLanguage()
234 msgpack::MapDocNode Kern) { in emitKernelAttrs() argument
237 Kern[".reqd_workgroup_size"] = getWorkGroupDimensions(Node); in emitKernelAttrs()
239 Kern[".workgroup_size_hint"] = getWorkGroupDimensions(Node); in emitKernelAttrs()
241 Kern[".vec_type_hint"] = Kern.getDocument()->getNode( in emitKernelAttrs()
248 Kern[".device_enqueue_symbol"] = Kern.getDocument()->getNode( in emitKernelAttrs()
253 Kern[".kind"] = Kern.getDocument()->getNode("init"); in emitKernelAttrs()
255 Kern[".kind"] = Kern.getDocument()->getNode("fini"); in emitKernelAttrs()
259 msgpack::MapDocNode Kern) { in emitKernelArgs() argument
268 Kern[".args"] = Args; in emitKernelArgs()
466 auto Kern = HSAMetadataDoc->getMapNode(); in getHSAKernelProps() local
469 Kern[".kernarg_segment_size"] = Kern.getDocument()->getNode( in getHSAKernelProps()
471 Kern[".group_segment_fixed_size"] = in getHSAKernelProps()
472 Kern.getDocument()->getNode(ProgramInfo.LDSSize); in getHSAKernelProps()
473 DelayedExprs->assignDocNode(Kern[".private_segment_fixed_size"], in getHSAKernelProps()
476 DelayedExprs->assignDocNode(Kern[".uses_dynamic_stack"], in getHSAKernelProps()
482 Kern[".workgroup_processor_mode"] = in getHSAKernelProps()
483 Kern.getDocument()->getNode(ProgramInfo.WgpMode); in getHSAKernelProps()
486 Kern[".kernarg_segment_align"] = in getHSAKernelProps()
487 Kern.getDocument()->getNode(std::max(Align(4), MaxKernArgAlign).value()); in getHSAKernelProps()
488 Kern[".wavefront_size"] = in getHSAKernelProps()
489 Kern.getDocument()->getNode(STM.getWavefrontSize()); in getHSAKernelProps()
490 DelayedExprs->assignDocNode(Kern[".sgpr_count"], msgpack::Type::UInt, in getHSAKernelProps()
492 DelayedExprs->assignDocNode(Kern[".vgpr_count"], msgpack::Type::UInt, in getHSAKernelProps()
497 DelayedExprs->assignDocNode(Kern[".agpr_count"], msgpack::Type::UInt, in getHSAKernelProps()
501 Kern[".max_flat_workgroup_size"] = in getHSAKernelProps()
502 Kern.getDocument()->getNode(MFI.getMaxFlatWorkGroupSize()); in getHSAKernelProps()
507 Kern[".max_num_workgroups_x"] = Kern.getDocument()->getNode(NumWGX); in getHSAKernelProps()
508 Kern[".max_num_workgroups_y"] = Kern.getDocument()->getNode(NumWGY); in getHSAKernelProps()
509 Kern[".max_num_workgroups_z"] = Kern.getDocument()->getNode(NumWGZ); in getHSAKernelProps()
511 Kern[".sgpr_spill_count"] = in getHSAKernelProps()
512 Kern.getDocument()->getNode(MFI.getNumSpilledSGPRs()); in getHSAKernelProps()
513 Kern[".vgpr_spill_count"] = in getHSAKernelProps()
514 Kern.getDocument()->getNode(MFI.getNumSpilledVGPRs()); in getHSAKernelProps()
516 return Kern; in getHSAKernelProps()
554 auto Kern = getHSAKernelProps(MF, ProgramInfo, CodeObjectVersion); in emitKernel() local
560 Kern[".name"] = Kern.getDocument()->getNode(Func.getName()); in emitKernel()
561 Kern[".symbol"] = Kern.getDocument()->getNode( in emitKernel()
563 emitKernelLanguage(Func, Kern); in emitKernel()
564 emitKernelAttrs(Func, Kern); in emitKernel()
565 emitKernelArgs(MF, Kern); in emitKernel()
568 Kernels.push_back(Kern); in emitKernel()
691 msgpack::MapDocNode Kern) { in emitKernelAttrs() argument
692 MetadataStreamerMsgPackV4::emitKernelAttrs(Func, Kern); in emitKernelAttrs()
695 Kern[".uniform_work_group_size"] = Kern.getDocument()->getNode(1); in emitKernelAttrs()