Changeset 4951 in CLRX


Ignore:
Timestamp:
Sep 7, 2019, 8:19:55 PM (2 months ago)
Author:
matszpk
Message:

CLRadeonExtender: ROCmBin: Add next stuff to generateROCmMetadataMsgPack.

File:
1 edited

Legend:

Unmodified
Added
Removed
  • CLRadeonExtender/trunk/amdbin/ROCmMetadataMP.cpp

    r4950 r4951  
    10811081    MsgPackArrayWriter putArray(size_t aelemsNum);
    10821082    MsgPackMapWriter putMap(size_t melemsNum);
    1083     void flush();
    10841083};
    10851084
     
    11001099    MsgPackMapWriter putValueMap(size_t melemsNum);
    11011100    std::vector<cxbyte>& putValueElement();
    1102     void flush();
    11031101};
    11041102
     
    12631261};
    12641262
    1265 
     1263// helper for checking whether value is supplied
     1264static inline bool hasValue(cxuint value)
     1265{ return value!=BINGEN_NOTSUPPLIED && value!=BINGEN_DEFAULT; }
     1266
     1267static inline bool hasValue(uint64_t value)
     1268{ return value!=BINGEN64_NOTSUPPLIED && value!=BINGEN64_DEFAULT; }
    12661269
    12671270void CLRX::generateROCmMetadataMsgPack(const ROCmMetadata& mdInfo,
     
    12731276    {
    12741277        const ROCmKernelMetadata& kernelMD = mdInfo.kernels[i];
    1275         const size_t mapSize = 15 + (!kernelMD.deviceEnqueueSymbol.empty()) +
     1278        const size_t mapSize = 13 + (!kernelMD.deviceEnqueueSymbol.empty()) +
    12761279                (kernelMD.reqdWorkGroupSize[0]!=0 ||
    12771280                 kernelMD.reqdWorkGroupSize[1]!=0 ||
     
    12801283                (kernelMD.workGroupSizeHint[0]!=0 ||
    12811284                 kernelMD.workGroupSizeHint[1]!=0 ||
    1282                  kernelMD.workGroupSizeHint[2]!=0);
     1285                 kernelMD.workGroupSizeHint[2]!=0) +
     1286                (!kernelMD.language.empty()) +
     1287                (kernelMD.langVersion[0]!=BINGEN_NOTSUPPLIED);
    12831288        MsgPackMapWriter kwriter = kernelsWriter.putMap(mapSize);
    12841289        kwriter.putKeyString(".args");
     
    13801385        }
    13811386        } //
    1382     }
    1383 }
     1387        if (!kernelMD.deviceEnqueueSymbol.empty())
     1388        {
     1389            kwriter.putKeyString(".device_enqueue_symbol");
     1390            kwriter.putValueString(kernelMD.deviceEnqueueSymbol.c_str());
     1391        }
     1392       
     1393        const ROCmKernelDescriptor& kdesc = *(kdescs[i]);
     1394       
     1395        kwriter.putKeyString(".group_segment_fixed_size");
     1396        kwriter.putValueUInt(hasValue(kernelMD.groupSegmentFixedSize) ?
     1397                kernelMD.groupSegmentFixedSize : kdesc.groupSegmentFixedSize);
     1398        kwriter.putKeyString(".kernarg_segment_align");
     1399        kwriter.putValueUInt(kernelMD.kernargSegmentAlign);
     1400        kwriter.putKeyString(".kernarg_segment_size");
     1401        kwriter.putValueUInt(kernelMD.kernargSegmentSize);
     1402       
     1403        if (!kernelMD.language.empty())
     1404        {
     1405            kwriter.putKeyString(".language");
     1406            kwriter.putValueString(kernelMD.language.c_str());
     1407        }
     1408        if (kernelMD.langVersion[0]!=BINGEN_NOTSUPPLIED)
     1409        {
     1410            kwriter.putKeyString(".language_version");
     1411            MsgPackArrayWriter verWriter = kwriter.putValueArray(2);
     1412            verWriter.putUInt(kernelMD.langVersion[0]);
     1413            verWriter.putUInt(kernelMD.langVersion[1]);
     1414        }
     1415       
     1416        kwriter.putKeyString(".max_flat_workgroup_size");
     1417        kwriter.putValueUInt(kernelMD.maxFlatWorkGroupSize);
     1418        kwriter.putKeyString(".name");
     1419        kwriter.putValueString(kernelMD.name.c_str());
     1420        kwriter.putKeyString(".private_segment_fixed_size");
     1421        kwriter.putValueUInt(hasValue(kernelMD.privateSegmentFixedSize) ?
     1422                kernelMD.privateSegmentFixedSize : kdesc.privateSegmentFixedSize);
     1423       
     1424        if (kernelMD.reqdWorkGroupSize[0] != 0 || kernelMD.reqdWorkGroupSize[1] != 0 ||
     1425            kernelMD.reqdWorkGroupSize[2] != 0)
     1426        {
     1427            kwriter.putKeyString(".reqd_workgroup_size");
     1428            MsgPackArrayWriter rwriter = kwriter.putValueArray(3);
     1429            for (cxuint i = 0; i < 3; i++)
     1430                rwriter.putUInt(kernelMD.reqdWorkGroupSize[i]);
     1431        }
     1432       
     1433        kwriter.putKeyString(".sgpr_count");
     1434        kwriter.putValueUInt(kernelMD.sgprsNum);
     1435        kwriter.putKeyString(".sgpr_spill_count");
     1436        kwriter.putValueUInt(kernelMD.spilledSgprs);
     1437        if (!kernelMD.vecTypeHint.empty())
     1438        {
     1439            kwriter.putKeyString(".vec_type_hint");
     1440            kwriter.putValueString(kernelMD.vecTypeHint.c_str());
     1441        }
     1442        kwriter.putKeyString(".vgpr_count");
     1443        kwriter.putValueUInt(kernelMD.vgprsNum);
     1444        kwriter.putKeyString(".vgpr_spill_count");
     1445        kwriter.putValueUInt(kernelMD.spilledVgprs);
     1446        kwriter.putKeyString(".wavefront_size");
     1447        kwriter.putValueUInt(kernelMD.wavefrontSize);
     1448       
     1449        if (kernelMD.workGroupSizeHint[0] != 0 || kernelMD.workGroupSizeHint[1] != 0 ||
     1450            kernelMD.workGroupSizeHint[2] != 0)
     1451        {
     1452            kwriter.putKeyString(".workgroup_size_hint");
     1453            MsgPackArrayWriter rwriter = kwriter.putValueArray(3);
     1454            for (cxuint i = 0; i < 3; i++)
     1455                rwriter.putUInt(kernelMD.workGroupSizeHint[i]);
     1456        }
     1457    }
     1458}
Note: See TracChangeset for help on using the changeset viewer.