Lines Matching defs:Kern
215 msgpack::MapDocNode Kern) {
224 Kern[".language"] = Kern.getDocument()->getNode("OpenCL C");
225 auto LanguageVersion = Kern.getDocument()->getArrayNode();
226 LanguageVersion.push_back(Kern.getDocument()->getNode(
228 LanguageVersion.push_back(Kern.getDocument()->getNode(
230 Kern[".language_version"] = LanguageVersion;
234 msgpack::MapDocNode Kern) {
237 Kern[".reqd_workgroup_size"] = getWorkGroupDimensions(Node);
239 Kern[".workgroup_size_hint"] = getWorkGroupDimensions(Node);
241 Kern[".vec_type_hint"] = Kern.getDocument()->getNode(
248 Kern[".device_enqueue_symbol"] = Kern.getDocument()->getNode(
253 Kern[".kind"] = Kern.getDocument()->getNode("init");
255 Kern[".kind"] = Kern.getDocument()->getNode("fini");
259 msgpack::MapDocNode Kern) {
272 Kern[".args"] = Args;
470 auto Kern = HSAMetadataDoc->getMapNode();
473 Kern[".kernarg_segment_size"] = Kern.getDocument()->getNode(
475 Kern[".group_segment_fixed_size"] =
476 Kern.getDocument()->getNode(ProgramInfo.LDSSize);
477 DelayedExprs->assignDocNode(Kern[".private_segment_fixed_size"],
480 DelayedExprs->assignDocNode(Kern[".uses_dynamic_stack"],
486 Kern[".workgroup_processor_mode"] =
487 Kern.getDocument()->getNode(ProgramInfo.WgpMode);
490 Kern[".kernarg_segment_align"] =
491 Kern.getDocument()->getNode(std::max(Align(4), MaxKernArgAlign).value());
492 Kern[".wavefront_size"] =
493 Kern.getDocument()->getNode(STM.getWavefrontSize());
494 DelayedExprs->assignDocNode(Kern[".sgpr_count"], msgpack::Type::UInt,
496 DelayedExprs->assignDocNode(Kern[".vgpr_count"], msgpack::Type::UInt,
501 DelayedExprs->assignDocNode(Kern[".agpr_count"], msgpack::Type::UInt,
505 Kern[".max_flat_workgroup_size"] =
506 Kern.getDocument()->getNode(MFI.getMaxFlatWorkGroupSize());
514 Kern[".max_num_workgroups_x"] = Kern.getDocument()->getNode(NumWGX);
517 Kern[".max_num_workgroups_y"] = Kern.getDocument()->getNode(NumWGY);
520 Kern[".max_num_workgroups_z"] = Kern.getDocument()->getNode(NumWGZ);
522 Kern[".sgpr_spill_count"] =
523 Kern.getDocument()->getNode(MFI.getNumSpilledSGPRs());
524 Kern[".vgpr_spill_count"] =
525 Kern.getDocument()->getNode(MFI.getNumSpilledVGPRs());
527 return Kern;
565 auto Kern = getHSAKernelProps(MF, ProgramInfo, CodeObjectVersion);
571 Kern[".name"] = Kern.getDocument()->getNode(Func.getName());
572 Kern[".symbol"] = Kern.getDocument()->getNode(
574 emitKernelLanguage(Func, Kern);
575 emitKernelAttrs(Func, Kern);
576 emitKernelArgs(MF, Kern);
579 Kernels.push_back(Kern);
702 msgpack::MapDocNode Kern) {
703 MetadataStreamerMsgPackV4::emitKernelAttrs(Func, Kern);
706 Kern[".uniform_work_group_size"] = Kern.getDocument()->getNode(1);