Wenkai Du
89349f2ce4
* Template unroll for RCCL kernels * Adding unroll template arg during CMake hipification * Reduce linking parallel jobs to avoid OOM in CI * Workaround issues with UT tests SWDEV-469533: register spill fix is needed for mainline build LWPCOMMLIBS-369: cannot enable 112 channels with 80 CUs Use -parallel-jobs=8 for linking * CI: do not use -j 16 when building * CI: use -j 8 when building * Only reduce parallel linking job for CI extended * Restore original jenkins command. Change parallel linking jobs in cmake * Disable MSCCLPP --------- Co-authored-by: gilbertlee-amd <gilbert.lee@amd.com>
107 wiersze
3.4 KiB
C++
107 wiersze
3.4 KiB
C++
#ifndef STANDALONE_UTILS_H
|
|
#define STANDALONE_UTILS_H
|
|
|
|
#include <iostream>
|
|
#include <cstdio>
|
|
#include <regex>
|
|
|
|
#define HIPCALL(cmd) \
|
|
do { \
|
|
hipError_t error = (cmd); \
|
|
if (error != hipSuccess) \
|
|
{ \
|
|
printf("Encountered HIP error (%s) at line %d in file %s\n", \
|
|
hipGetErrorString(error), __LINE__, __FILE__); \
|
|
exit(-1); \
|
|
} \
|
|
} while (0)
|
|
|
|
#define NCCLCHECK(cmd) do { \
|
|
ncclResult_t res = cmd; \
|
|
if (res != ncclSuccess) { \
|
|
printf("NCCL failure %s:%d '%s'\n", \
|
|
__FILE__,__LINE__,ncclGetErrorString(res)); \
|
|
} \
|
|
} while(0)
|
|
|
|
// should be 112, temp fix to make CI pass
|
|
#define MAX_STACK_SIZE 360
|
|
|
|
#ifdef ENABLE_LL128
|
|
#define MAX_STACK_SIZE_gfx90a 296
|
|
#else
|
|
#define MAX_STACK_SIZE_gfx90a MAX_STACK_SIZE
|
|
#endif
|
|
|
|
struct KernelInfo {
|
|
std::string name;
|
|
int privateSegmentFixedSize = 0;
|
|
};
|
|
|
|
struct ArchInfo {
|
|
std::string archName;
|
|
std::vector<KernelInfo> kernels;
|
|
};
|
|
|
|
std::string executeCommand(const char* cmd) {
|
|
std::string result;
|
|
FILE* pipe = popen(cmd, "r");
|
|
|
|
if (!pipe) {
|
|
std::cerr << "Error executing command: " << cmd << std::endl;
|
|
return result;
|
|
}
|
|
|
|
char buffer[128];
|
|
while (!feof(pipe)) {
|
|
if (fgets(buffer, 128, pipe) != NULL) {
|
|
result += buffer;
|
|
}
|
|
}
|
|
|
|
pclose(pipe);
|
|
return result;
|
|
}
|
|
|
|
std::vector<std::string> splitString(const std::string& str, char delimiter) {
|
|
std::vector<std::string> result;
|
|
std::istringstream iss(str);
|
|
|
|
std::string line;
|
|
while(std::getline(iss, line, delimiter)) {
|
|
result.push_back(line);
|
|
}
|
|
|
|
return result;
|
|
}
|
|
|
|
|
|
ArchInfo parseMetadata(const std::vector<std::string>& list) {
|
|
ArchInfo archInfo;
|
|
KernelInfo currKernelInfo;
|
|
|
|
std::regex amdhsaTargetRegex("amdhsa.target:\\s+(?:'?)amdgcn-amd-amdhsa--(\\w+)(?:'?)");
|
|
std::regex kernelNameRegex("\\.name:\\s+(\\w+)");
|
|
std::regex privateSegmentSizeRegex("\\.private_segment_fixed_size:\\s+(\\d+)");
|
|
|
|
for (const auto& line : list) {
|
|
std::smatch match;
|
|
|
|
if (std::regex_search(line, match, amdhsaTargetRegex)) {
|
|
archInfo.archName = match[1];
|
|
} else if (std::regex_search(line, match, kernelNameRegex)) {
|
|
currKernelInfo.name = match[1];
|
|
} else if (std::regex_search(line, match, privateSegmentSizeRegex)) {
|
|
currKernelInfo.privateSegmentFixedSize = std::stoi(match[1]);
|
|
}
|
|
|
|
if (!currKernelInfo.name.empty() && currKernelInfo.privateSegmentFixedSize != 0) {
|
|
archInfo.kernels.push_back(currKernelInfo);
|
|
currKernelInfo = {}; // Empty kernelInfo
|
|
}
|
|
}
|
|
|
|
return archInfo;
|
|
}
|
|
|
|
#endif |