15#define GPUCA_GPUCODE_HOSTONLY
19#include "GPUDefParametersLoad.inc"
26#include <oneapi/tbb.h>
35#include "GPUNoFastMathKernels.h"
37int32_t GPUReconstructionCUDA::genRTC(std::string&
filename, uint32_t& nCompile)
39 std::string rtcparam = std::string(
"#define GPUCA_RTC_CODE\n") +
40 std::string(
GetProcessingSettings().rtc.optSpecialCode ?
"#define GPUCA_RTC_SPECIAL_CODE(...) __VA_ARGS__\n" :
"#define GPUCA_RTC_SPECIAL_CODE(...)\n") +
50 std::vector<std::string> kernels;
51 getRTCKernelCalls(kernels);
52 std::string kernelsall;
53 for (uint32_t
i = 0;
i < kernels.size();
i++) {
54 kernelsall += kernels[
i] +
"\n";
58 baseCommand += (getenv(
"O2_GPU_RTC_OVERRIDE_CMD") ? std::string(getenv(
"O2_GPU_RTC_OVERRIDE_CMD")) :
std::
string(_binary_GPUReconstructionCUDArtc_command_start, _binary_GPUReconstructionCUDArtc_command_len));
59 baseCommand += std::string(
" ") + (
GetProcessingSettings().rtctech.overrideArchitecture !=
"" ?
GetProcessingSettings().rtctech.overrideArchitecture : std::string(_binary_GPUReconstructionCUDArtc_command_arch_start, _binary_GPUReconstructionCUDArtc_command_arch_len));
64 throw std::runtime_error(
"Cannot open launch bounds parameter module file");
66 fseek(fp, 0, SEEK_END);
67 size_t size = ftell(fp);
69 throw std::runtime_error(
"launch bounds parameter file has incorrect size");
71 fseek(fp, 0, SEEK_SET);
73 throw std::runtime_error(
"Error reading launch bounds parameter file");
77 const std::string launchBounds = o2::gpu::internal::GPUDefParametersExport(*
mParDevice,
true) +
80 GPUInfo(
"RTC Launch Bounds:\n%s", launchBounds.c_str());
85 char shasource[21], shaparam[21], shacmd[21], shakernels[21], shabounds[21], shaversion[21];
87 o2::framework::internal::SHA1(shasource, _binary_GPUReconstructionCUDArtc_src_start, _binary_GPUReconstructionCUDArtc_src_len);
88 o2::framework::internal::SHA1(shaparam, rtcparam.c_str(), rtcparam.size());
89 o2::framework::internal::SHA1(shacmd, baseCommand.c_str(), baseCommand.size());
90 o2::framework::internal::SHA1(shakernels, kernelsall.c_str(), kernelsall.size());
91 o2::framework::internal::SHA1(shabounds, launchBounds.c_str(), launchBounds.size());
92 o2::framework::internal::SHA1(shaversion, compilerVersions.c_str(), compilerVersions.size());
96 bool cacheLoaded =
false;
103 mode_t
mask = S_IRUSR | S_IWUSR | S_IRGRP | S_IWGRP | S_IROTH | S_IWOTH;
106 throw std::runtime_error(
"Error opening rtc cache mutex lock file");
109 if (lockf(fd, F_LOCK, 0)) {
110 throw std::runtime_error(
"Error locking rtc cache mutex file");
119 auto checkSHA = [&](
const char* shacmp,
const char*
name) {
120 if (fread(sharead, 1, 20, fp) != 20) {
121 throw std::runtime_error(
"Cache file corrupt");
124 char shaprint1[41], shaprint2[41];
125 for (uint32_t
i = 0;
i < 20;
i++) {
126 sprintf(shaprint1 + 2 *
i,
"%02X ", shacmp[
i]);
127 sprintf(shaprint2 + 2 *
i,
"%02X ", sharead[
i]);
129 GPUInfo(
"SHA for %s: expected %s, read %s",
name, shaprint1, shaprint2);
132 GPUInfo(
"Cache file content outdated (%s)",
name);
137 if (checkSHA(shasource,
"source") ||
138 checkSHA(shaparam,
"param") ||
139 checkSHA(shacmd,
"command line") ||
140 checkSHA(shakernels,
"kernel definitions") ||
141 checkSHA(shabounds,
"launch bounds") ||
142 checkSHA(shaversion,
"compiler versions")) {
145 GPUSettingsProcessingRTC cachedSettings;
146 static_assert(std::is_trivially_copyable_v<GPUSettingsProcessingRTC> ==
true,
"GPUSettingsProcessingRTC must be POD");
147 if (fread(&cachedSettings,
sizeof(cachedSettings), 1, fp) != 1) {
148 throw std::runtime_error(
"Cache file corrupt");
151 GPUInfo(
"Cache file content outdated (rtc parameters)");
155 for (uint32_t
i = 0;
i < nCompile;
i++) {
156 if (fread(&
len,
sizeof(
len), 1, fp) != 1) {
157 throw std::runtime_error(
"Cache file corrupt");
161 throw std::runtime_error(
"Cache file corrupt");
164 if (fp2 ==
nullptr) {
165 throw std::runtime_error(
"Cannot open tmp file");
168 throw std::runtime_error(
"Error writing file");
172 GPUInfo(
"Using RTC cache file");
181 GPUInfo(
"Starting CUDA RTC Compilation");
185 tbb::parallel_for<uint32_t>(0, nCompile, [&](
auto i) {
191 throw std::runtime_error(
"Error opening file");
194 std::string kernel =
"extern \"C\" {";
199 const std::string deterministicStr = std::string(deterministic ?
"#define GPUCA_DETERMINISTIC_CODE(det, indet) det\n" :
"#define GPUCA_DETERMINISTIC_CODE(det, indet) indet\n");
201 if (fwrite(deterministicStr.c_str(), 1, deterministicStr.size(), fp) != deterministicStr.size() ||
202 fwrite(rtcparam.c_str(), 1, rtcparam.size(), fp) != rtcparam.size() ||
203 fwrite(launchBounds.c_str(), 1, launchBounds.size(), fp) != launchBounds.size() ||
204 fwrite(_binary_GPUReconstructionCUDArtc_src_start, 1, _binary_GPUReconstructionCUDArtc_src_len, fp) != _binary_GPUReconstructionCUDArtc_src_len ||
205 fwrite(kernel.c_str(), 1, kernel.size(), fp) != kernel.size()) {
206 throw std::runtime_error(
"Error writing file");
209 std::string command = baseCommand;
211 command += std::string(
" ") + std::string(_binary_GPUReconstructionCUDArtc_command_no_fast_math_start, _binary_GPUReconstructionCUDArtc_command_no_fast_math_len);
215 command +=
" &> /dev/null";
217 command +=
" > /dev/null";
220 printf(
"Running command %s\n", command.c_str());
222 if (system(command.c_str())) {
224 printf(
"Source code file: %s",
filename.c_str());
226 throw std::runtime_error(
"Error during CUDA compilation");
228 }, tbb::simple_partitioner());
235 throw std::runtime_error(
"Cannot open cache file for writing");
237 GPUInfo(
"Storing RTC compilation result in cache file");
239 if (fwrite(shasource, 1, 20, fp) != 20 ||
240 fwrite(shaparam, 1, 20, fp) != 20 ||
241 fwrite(shacmd, 1, 20, fp) != 20 ||
242 fwrite(shakernels, 1, 20, fp) != 20 ||
243 fwrite(shabounds, 1, 20, fp) != 20 ||
244 fwrite(shaversion, 1, 20, fp) != 20 ||
246 throw std::runtime_error(
"Error writing cache file");
250 for (uint32_t
i = 0;
i < nCompile;
i++) {
252 if (fp2 ==
nullptr) {
253 throw std::runtime_error(
"Cannot open cuda module file");
255 fseek(fp2, 0, SEEK_END);
256 size_t size = ftell(fp2);
258 fseek(fp2, 0, SEEK_SET);
260 throw std::runtime_error(
"Error reading cuda module file");
264 if (fwrite(&
size,
sizeof(
size), 1, fp) != 1 ||
266 throw std::runtime_error(
"Error writing cache file");
273 if (lockf(fd, F_ULOCK, 0)) {
274 throw std::runtime_error(
"Error unlocking RTC cache mutex file");
double GetCurrentElapsedTime(bool reset=false)
static const char * GetKernelName()
GPUDefParameters * mParDevice
static std::string getBackendVersions()
const GPUSettingsProcessing & GetProcessingSettings() const
GLsizei const GLchar *const * string
GLuint const GLchar * name
GLenum GLenum GLsizei len
Polygon< T > close(Polygon< T > polygon)
Defining DataPointCompositeObject explicitly as copiable.
std::string to_string(gsl::span< T, Size > span)
#define QGET_LD_BINARY_SYMBOLS(filename)