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") +
49 std::vector<std::string> kernels;
50 getRTCKernelCalls(kernels);
51 std::string kernelsall;
52 for (uint32_t
i = 0;
i < kernels.size();
i++) {
53 kernelsall += kernels[
i] +
"\n";
57 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));
58 baseCommand += std::string(
" ") + (
GetProcessingSettings().rtctech.overrideArchitecture !=
"" ?
GetProcessingSettings().rtctech.overrideArchitecture : std::string(_binary_GPUReconstructionCUDArtc_command_arch_start, _binary_GPUReconstructionCUDArtc_command_arch_len));
63 throw std::runtime_error(
"Cannot open launch bounds parameter module file");
65 fseek(fp, 0, SEEK_END);
66 size_t size = ftell(fp);
68 throw std::runtime_error(
"launch bounds parameter file has incorrect size");
70 fseek(fp, 0, SEEK_SET);
72 throw std::runtime_error(
"Error reading launch bounds parameter file");
76 const std::string launchBounds = o2::gpu::internal::GPUDefParametersExport(*
mParDevice,
true) +
79 GPUInfo(
"RTC Launch Bounds:\n%s", launchBounds.c_str());
82 char shasource[21], shaparam[21], shacmd[21], shakernels[21], shabounds[21];
84 o2::framework::internal::SHA1(shasource, _binary_GPUReconstructionCUDArtc_src_start, _binary_GPUReconstructionCUDArtc_src_len);
85 o2::framework::internal::SHA1(shaparam, rtcparam.c_str(), rtcparam.size());
86 o2::framework::internal::SHA1(shacmd, baseCommand.c_str(), baseCommand.size());
87 o2::framework::internal::SHA1(shakernels, kernelsall.c_str(), kernelsall.size());
88 o2::framework::internal::SHA1(shabounds, launchBounds.c_str(), launchBounds.size());
92 bool cacheLoaded =
false;
99 mode_t
mask = S_IRUSR | S_IWUSR | S_IRGRP | S_IWGRP | S_IROTH | S_IWOTH;
102 throw std::runtime_error(
"Error opening rtc cache mutex lock file");
105 if (lockf(fd, F_LOCK, 0)) {
106 throw std::runtime_error(
"Error locking rtc cache mutex file");
115 auto checkSHA = [&](
const char* shacmp,
const char*
name) {
116 if (fread(sharead, 1, 20, fp) != 20) {
117 throw std::runtime_error(
"Cache file corrupt");
120 char shaprint1[41], shaprint2[41];
121 for (uint32_t
i = 0;
i < 20;
i++) {
122 sprintf(shaprint1 + 2 *
i,
"%02X ", shacmp[
i]);
123 sprintf(shaprint2 + 2 *
i,
"%02X ", sharead[
i]);
125 GPUInfo(
"SHA for %s: expected %s, read %s",
name, shaprint1, shaprint2);
128 GPUInfo(
"Cache file content outdated (%s)",
name);
133 if (checkSHA(shasource,
"source") ||
134 checkSHA(shaparam,
"param") ||
135 checkSHA(shacmd,
"command line") ||
136 checkSHA(shakernels,
"kernel definitions") ||
137 checkSHA(shabounds,
"launch bounds")) {
140 GPUSettingsProcessingRTC cachedSettings;
141 static_assert(std::is_trivially_copyable_v<GPUSettingsProcessingRTC> ==
true,
"GPUSettingsProcessingRTC must be POD");
142 if (fread(&cachedSettings,
sizeof(cachedSettings), 1, fp) != 1) {
143 throw std::runtime_error(
"Cache file corrupt");
146 GPUInfo(
"Cache file content outdated (rtc parameters)");
150 for (uint32_t
i = 0;
i < nCompile;
i++) {
151 if (fread(&
len,
sizeof(
len), 1, fp) != 1) {
152 throw std::runtime_error(
"Cache file corrupt");
156 throw std::runtime_error(
"Cache file corrupt");
159 if (fp2 ==
nullptr) {
160 throw std::runtime_error(
"Cannot open tmp file");
163 throw std::runtime_error(
"Error writing file");
167 GPUInfo(
"Using RTC cache file");
176 GPUInfo(
"Starting CUDA RTC Compilation");
180 tbb::parallel_for<uint32_t>(0, nCompile, [&](
auto i) {
186 throw std::runtime_error(
"Error opening file");
189 std::string kernel =
"extern \"C\" {";
194 const std::string deterministicStr = std::string(deterministic ?
"#define GPUCA_DETERMINISTIC_CODE(det, indet) det\n" :
"#define GPUCA_DETERMINISTIC_CODE(det, indet) indet\n");
196 if (fwrite(deterministicStr.c_str(), 1, deterministicStr.size(), fp) != deterministicStr.size() ||
197 fwrite(rtcparam.c_str(), 1, rtcparam.size(), fp) != rtcparam.size() ||
198 fwrite(launchBounds.c_str(), 1, launchBounds.size(), fp) != launchBounds.size() ||
199 fwrite(_binary_GPUReconstructionCUDArtc_src_start, 1, _binary_GPUReconstructionCUDArtc_src_len, fp) != _binary_GPUReconstructionCUDArtc_src_len ||
200 fwrite(kernel.c_str(), 1, kernel.size(), fp) != kernel.size()) {
201 throw std::runtime_error(
"Error writing file");
204 std::string command = baseCommand;
206 command += std::string(
" ") + std::string(_binary_GPUReconstructionCUDArtc_command_no_fast_math_start, _binary_GPUReconstructionCUDArtc_command_no_fast_math_len);
210 command +=
" &> /dev/null";
212 command +=
" > /dev/null";
215 printf(
"Running command %s\n", command.c_str());
217 if (system(command.c_str())) {
219 printf(
"Source code file: %s",
filename.c_str());
221 throw std::runtime_error(
"Error during CUDA compilation");
223 }, tbb::simple_partitioner());
230 throw std::runtime_error(
"Cannot open cache file for writing");
232 GPUInfo(
"Storing RTC compilation result in cache file");
234 if (fwrite(shasource, 1, 20, fp) != 20 ||
235 fwrite(shaparam, 1, 20, fp) != 20 ||
236 fwrite(shacmd, 1, 20, fp) != 20 ||
237 fwrite(shakernels, 1, 20, fp) != 20 ||
238 fwrite(shabounds, 1, 20, fp) != 20 ||
240 throw std::runtime_error(
"Error writing cache file");
244 for (uint32_t
i = 0;
i < nCompile;
i++) {
246 if (fp2 ==
nullptr) {
247 throw std::runtime_error(
"Cannot open cuda module file");
249 fseek(fp2, 0, SEEK_END);
250 size_t size = ftell(fp2);
252 fseek(fp2, 0, SEEK_SET);
254 throw std::runtime_error(
"Error reading cuda module file");
258 if (fwrite(&
size,
sizeof(
size), 1, fp) != 1 ||
260 throw std::runtime_error(
"Error writing cache file");
267 if (lockf(fd, F_ULOCK, 0)) {
268 throw std::runtime_error(
"Error unlocking RTC cache mutex file");
double GetCurrentElapsedTime(bool reset=false)
static const char * GetKernelName()
GPUDefParameters * mParDevice
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)