Project
Loading...
Searching...
No Matches
GPUReconstructionCUDAGenRTC.cxx
Go to the documentation of this file.
1// Copyright 2019-2020 CERN and copyright holders of ALICE O2.
2// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders.
3// All rights not expressly granted are reserved.
4//
5// This software is distributed under the terms of the GNU General Public
6// License v3 (GPL Version 3), copied verbatim in the file "COPYING".
7//
8// In applying this license CERN does not waive the privileges and immunities
9// granted to it by virtue of its status as an Intergovernmental Organization
10// or submit itself to any jurisdiction.
11
14
15#define GPUCA_GPUCODE_HOSTONLY
16
18#include "GPUParamRTC.h"
19#include "GPUDefParametersLoad.inc"
20#include <unistd.h>
21#include "Framework/SHA1.h"
22#include <sys/stat.h>
23#include <fcntl.h>
24#include <filesystem>
25
26#include <oneapi/tbb.h>
27using namespace o2::gpu;
28
30QGET_LD_BINARY_SYMBOLS(GPUReconstructionCUDArtc_src);
31QGET_LD_BINARY_SYMBOLS(GPUReconstructionCUDArtc_command);
32QGET_LD_BINARY_SYMBOLS(GPUReconstructionCUDArtc_command_arch);
33QGET_LD_BINARY_SYMBOLS(GPUReconstructionCUDArtc_command_no_fast_math);
34
35#include "GPUNoFastMathKernels.h"
36
37int32_t GPUReconstructionCUDA::genRTC(std::string& filename, uint32_t& nCompile)
38{
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") +
41 std::string(GetProcessingSettings().rtc.optConstexpr ? "#define GPUCA_RTC_CONSTEXPR constexpr\n" : "#define GPUCA_RTC_CONSTEXPR\n") +
42 GPUParamRTC::generateRTCCode(param(), GetProcessingSettings().rtc.optConstexpr);
43 if (filename == "") {
44 filename = "/tmp/o2cagpu_rtc_";
45 }
46 filename += std::to_string(getpid());
47 filename += "_";
48 filename += std::to_string(rand());
49
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";
55 }
56
57 std::string baseCommand = (GetProcessingSettings().rtctech.prependCommand != "" ? (GetProcessingSettings().rtctech.prependCommand + " ") : "");
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));
60
61 if (GetProcessingSettings().rtctech.loadLaunchBoundsFromFile.size()) {
62 FILE* fp = fopen(GetProcessingSettings().rtctech.loadLaunchBoundsFromFile.c_str(), "rb");
63 if (fp == nullptr) {
64 throw std::runtime_error("Cannot open launch bounds parameter module file");
65 }
66 fseek(fp, 0, SEEK_END);
67 size_t size = ftell(fp);
68 if (size != sizeof(*mParDevice)) {
69 throw std::runtime_error("launch bounds parameter file has incorrect size");
70 }
71 fseek(fp, 0, SEEK_SET);
72 if (fread(mParDevice, 1, size, fp) != size) {
73 throw std::runtime_error("Error reading launch bounds parameter file");
74 }
75 fclose(fp);
76 }
77 const std::string launchBounds = o2::gpu::internal::GPUDefParametersExport(*mParDevice, true) +
78 "#define GPUCA_WARP_SIZE " + std::to_string(mWarpSize) + "\n";
79 if (GetProcessingSettings().rtctech.printLaunchBounds || GetProcessingSettings().debugLevel >= 3) {
80 GPUInfo("RTC Launch Bounds:\n%s", launchBounds.c_str());
81 }
82
83 const std::string compilerVersions = getBackendVersions();
84
85 char shasource[21], shaparam[21], shacmd[21], shakernels[21], shabounds[21], shaversion[21];
86 if (GetProcessingSettings().rtc.cacheOutput) {
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());
93 }
94
95 nCompile = GetProcessingSettings().rtc.compilePerKernel ? kernels.size() : 1;
96 bool cacheLoaded = false;
97 int32_t fd = 0;
98 if (GetProcessingSettings().rtc.cacheOutput) {
99 if (GetProcessingSettings().rtctech.cacheFolder != ".") {
100 std::filesystem::create_directories(GetProcessingSettings().rtctech.cacheFolder);
101 }
102 if (GetProcessingSettings().rtctech.cacheMutex) {
103 mode_t mask = S_IRUSR | S_IWUSR | S_IRGRP | S_IWGRP | S_IROTH | S_IWOTH;
104 fd = open((GetProcessingSettings().rtctech.cacheFolder + "/cache.lock").c_str(), O_RDWR | O_CREAT | O_CLOEXEC, mask);
105 if (fd == -1) {
106 throw std::runtime_error("Error opening rtc cache mutex lock file");
107 }
108 fchmod(fd, mask);
109 if (lockf(fd, F_LOCK, 0)) {
110 throw std::runtime_error("Error locking rtc cache mutex file");
111 }
112 }
113
114 FILE* fp = fopen((GetProcessingSettings().rtctech.cacheFolder + "/rtc.cuda.cache").c_str(), "rb");
115 char sharead[20];
116 if (fp) {
117 size_t len;
118 while (true) {
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");
122 }
123 if (GetProcessingSettings().debugLevel >= 3) {
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]);
128 }
129 GPUInfo("SHA for %s: expected %s, read %s", name, shaprint1, shaprint2);
130 }
131 if (!GetProcessingSettings().rtctech.ignoreCacheValid && memcmp(sharead, shacmp, 20)) {
132 GPUInfo("Cache file content outdated (%s)", name);
133 return 1;
134 }
135 return 0;
136 };
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")) {
143 break;
144 }
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");
149 }
150 if (!GetProcessingSettings().rtctech.ignoreCacheValid && !(cachedSettings == GetProcessingSettings().rtc)) {
151 GPUInfo("Cache file content outdated (rtc parameters)");
152 break;
153 }
154 std::vector<char> buffer;
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");
158 }
159 buffer.resize(len);
160 if (fread(buffer.data(), 1, len, fp) != len) {
161 throw std::runtime_error("Cache file corrupt");
162 }
163 FILE* fp2 = fopen((filename + "_" + std::to_string(i) + mRtcBinExtension).c_str(), "w+b");
164 if (fp2 == nullptr) {
165 throw std::runtime_error("Cannot open tmp file");
166 }
167 if (fwrite(buffer.data(), 1, len, fp2) != len) {
168 throw std::runtime_error("Error writing file");
169 }
170 fclose(fp2);
171 }
172 GPUInfo("Using RTC cache file");
173 cacheLoaded = true;
174 break;
175 };
176 fclose(fp);
177 }
178 }
179 if (!cacheLoaded) {
180 if (GetProcessingSettings().debugLevel >= 0) {
181 GPUInfo("Starting CUDA RTC Compilation");
182 }
183 HighResTimer rtcTimer;
184 rtcTimer.ResetStart();
185 tbb::parallel_for<uint32_t>(0, nCompile, [&](auto i) {
186 if (GetProcessingSettings().debugLevel >= 3) {
187 printf("Compiling %s\n", (filename + "_" + std::to_string(i) + mRtcSrcExtension).c_str());
188 }
189 FILE* fp = fopen((filename + "_" + std::to_string(i) + mRtcSrcExtension).c_str(), "w+b");
190 if (fp == nullptr) {
191 throw std::runtime_error("Error opening file");
192 }
193
194 std::string kernel = "extern \"C\" {";
195 kernel += GetProcessingSettings().rtc.compilePerKernel ? kernels[i] : kernelsall;
196 kernel += "}";
197
198 bool deterministic = GetProcessingSettings().rtc.deterministic || (GetProcessingSettings().rtc.compilePerKernel && o2::gpu::internal::noFastMathKernels.find(GetKernelName(i)) != o2::gpu::internal::noFastMathKernels.end());
199 const std::string deterministicStr = std::string(deterministic ? "#define GPUCA_DETERMINISTIC_CODE(det, indet) det\n" : "#define GPUCA_DETERMINISTIC_CODE(det, indet) indet\n");
200
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");
207 }
208 fclose(fp);
209 std::string command = baseCommand;
210 if (deterministic) {
211 command += std::string(" ") + std::string(_binary_GPUReconstructionCUDArtc_command_no_fast_math_start, _binary_GPUReconstructionCUDArtc_command_no_fast_math_len);
212 }
213 command += " -c " + filename + "_" + std::to_string(i) + mRtcSrcExtension + " -o " + filename + "_" + std::to_string(i) + mRtcBinExtension;
214 if (GetProcessingSettings().debugLevel < 0) {
215 command += " &> /dev/null";
216 } else if (GetProcessingSettings().debugLevel < 2) {
217 command += " > /dev/null";
218 }
219 if (GetProcessingSettings().debugLevel >= 3) {
220 printf("Running command %s\n", command.c_str());
221 }
222 if (system(command.c_str())) {
223 if (GetProcessingSettings().debugLevel >= 3) {
224 printf("Source code file: %s", filename.c_str());
225 }
226 throw std::runtime_error("Error during CUDA compilation");
227 } // clang-format off
228 }, tbb::simple_partitioner()); // clang-format on
229 if (GetProcessingSettings().debugLevel >= 0) {
230 GPUInfo("RTC Compilation finished (%f seconds)", rtcTimer.GetCurrentElapsedTime());
231 }
232 if (GetProcessingSettings().rtc.cacheOutput) {
233 FILE* fp = fopen((GetProcessingSettings().rtctech.cacheFolder + "/rtc.cuda.cache").c_str(), "w+b");
234 if (fp == nullptr) {
235 throw std::runtime_error("Cannot open cache file for writing");
236 }
237 GPUInfo("Storing RTC compilation result in cache file");
238
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 ||
245 fwrite(&GetProcessingSettings().rtc, sizeof(GetProcessingSettings().rtc), 1, fp) != 1) {
246 throw std::runtime_error("Error writing cache file");
247 }
248
249 std::vector<char> buffer;
250 for (uint32_t i = 0; i < nCompile; i++) {
251 FILE* fp2 = fopen((filename + "_" + std::to_string(i) + mRtcBinExtension).c_str(), "rb");
252 if (fp2 == nullptr) {
253 throw std::runtime_error("Cannot open cuda module file");
254 }
255 fseek(fp2, 0, SEEK_END);
256 size_t size = ftell(fp2);
257 buffer.resize(size);
258 fseek(fp2, 0, SEEK_SET);
259 if (fread(buffer.data(), 1, size, fp2) != size) {
260 throw std::runtime_error("Error reading cuda module file");
261 }
262 fclose(fp2);
263
264 if (fwrite(&size, sizeof(size), 1, fp) != 1 ||
265 fwrite(buffer.data(), 1, size, fp) != size) {
266 throw std::runtime_error("Error writing cache file");
267 }
268 }
269 fclose(fp);
270 }
271 }
272 if (GetProcessingSettings().rtc.cacheOutput && GetProcessingSettings().rtctech.cacheMutex) {
273 if (lockf(fd, F_ULOCK, 0)) {
274 throw std::runtime_error("Error unlocking RTC cache mutex file");
275 }
276 close(fd);
277 }
278
279 return 0;
280}
int32_t i
double GetCurrentElapsedTime(bool reset=false)
Definition timer.cxx:110
void ResetStart()
Definition timer.cxx:63
static const char * GetKernelName()
const GPUSettingsProcessing & GetProcessingSettings() const
GLsizei const GLchar *const * string
Definition glcorearb.h:809
GLuint buffer
Definition glcorearb.h:655
GLsizeiptr size
Definition glcorearb.h:659
GLuint const GLchar * name
Definition glcorearb.h:781
GLenum GLenum GLsizei len
Definition glcorearb.h:4232
GLenum GLfloat param
Definition glcorearb.h:271
GLint GLuint mask
Definition glcorearb.h:291
Polygon< T > close(Polygon< T > polygon)
Definition Polygon.h:126
Defining DataPointCompositeObject explicitly as copiable.
std::string to_string(gsl::span< T, Size > span)
Definition common.h:52
std::string filename()
#define QGET_LD_BINARY_SYMBOLS(filename)