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 GPUParamRTC::generateRTCCode(param(), GetProcessingSettings().rtc.optConstexpr);
42 if (filename == "") {
43 filename = "/tmp/o2cagpu_rtc_";
44 }
45 filename += std::to_string(getpid());
46 filename += "_";
47 filename += std::to_string(rand());
48
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";
54 }
55
56 std::string baseCommand = (GetProcessingSettings().rtctech.prependCommand != "" ? (GetProcessingSettings().rtctech.prependCommand + " ") : "");
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));
59
60 if (GetProcessingSettings().rtctech.loadLaunchBoundsFromFile.size()) {
61 FILE* fp = fopen(GetProcessingSettings().rtctech.loadLaunchBoundsFromFile.c_str(), "rb");
62 if (fp == nullptr) {
63 throw std::runtime_error("Cannot open launch bounds parameter module file");
64 }
65 fseek(fp, 0, SEEK_END);
66 size_t size = ftell(fp);
67 if (size != sizeof(*mParDevice)) {
68 throw std::runtime_error("launch bounds parameter file has incorrect size");
69 }
70 fseek(fp, 0, SEEK_SET);
71 if (fread(mParDevice, 1, size, fp) != size) {
72 throw std::runtime_error("Error reading launch bounds parameter file");
73 }
74 fclose(fp);
75 }
76 const std::string launchBounds = o2::gpu::internal::GPUDefParametersExport(*mParDevice, true) +
77 "#define GPUCA_WARP_SIZE " + std::to_string(mWarpSize) + "\n";
78 if (GetProcessingSettings().rtctech.printLaunchBounds || GetProcessingSettings().debugLevel >= 3) {
79 GPUInfo("RTC Launch Bounds:\n%s", launchBounds.c_str());
80 }
81
82 char shasource[21], shaparam[21], shacmd[21], shakernels[21], shabounds[21];
83 if (GetProcessingSettings().rtc.cacheOutput) {
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());
89 }
90
91 nCompile = GetProcessingSettings().rtc.compilePerKernel ? kernels.size() : 1;
92 bool cacheLoaded = false;
93 int32_t fd = 0;
94 if (GetProcessingSettings().rtc.cacheOutput) {
95 if (GetProcessingSettings().rtctech.cacheFolder != ".") {
96 std::filesystem::create_directories(GetProcessingSettings().rtctech.cacheFolder);
97 }
98 if (GetProcessingSettings().rtctech.cacheMutex) {
99 mode_t mask = S_IRUSR | S_IWUSR | S_IRGRP | S_IWGRP | S_IROTH | S_IWOTH;
100 fd = open((GetProcessingSettings().rtctech.cacheFolder + "/cache.lock").c_str(), O_RDWR | O_CREAT | O_CLOEXEC, mask);
101 if (fd == -1) {
102 throw std::runtime_error("Error opening rtc cache mutex lock file");
103 }
104 fchmod(fd, mask);
105 if (lockf(fd, F_LOCK, 0)) {
106 throw std::runtime_error("Error locking rtc cache mutex file");
107 }
108 }
109
110 FILE* fp = fopen((GetProcessingSettings().rtctech.cacheFolder + "/rtc.cuda.cache").c_str(), "rb");
111 char sharead[20];
112 if (fp) {
113 size_t len;
114 while (true) {
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");
118 }
119 if (GetProcessingSettings().debugLevel >= 3) {
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]);
124 }
125 GPUInfo("SHA for %s: expected %s, read %s", name, shaprint1, shaprint2);
126 }
127 if (!GetProcessingSettings().rtctech.ignoreCacheValid && memcmp(sharead, shacmp, 20)) {
128 GPUInfo("Cache file content outdated (%s)", name);
129 return 1;
130 }
131 return 0;
132 };
133 if (checkSHA(shasource, "source") ||
134 checkSHA(shaparam, "param") ||
135 checkSHA(shacmd, "command line") ||
136 checkSHA(shakernels, "kernel definitions") ||
137 checkSHA(shabounds, "launch bounds")) {
138 break;
139 }
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");
144 }
145 if (!GetProcessingSettings().rtctech.ignoreCacheValid && !(cachedSettings == GetProcessingSettings().rtc)) {
146 GPUInfo("Cache file content outdated (rtc parameters)");
147 break;
148 }
149 std::vector<char> buffer;
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");
153 }
154 buffer.resize(len);
155 if (fread(buffer.data(), 1, len, fp) != len) {
156 throw std::runtime_error("Cache file corrupt");
157 }
158 FILE* fp2 = fopen((filename + "_" + std::to_string(i) + mRtcBinExtension).c_str(), "w+b");
159 if (fp2 == nullptr) {
160 throw std::runtime_error("Cannot open tmp file");
161 }
162 if (fwrite(buffer.data(), 1, len, fp2) != len) {
163 throw std::runtime_error("Error writing file");
164 }
165 fclose(fp2);
166 }
167 GPUInfo("Using RTC cache file");
168 cacheLoaded = true;
169 break;
170 };
171 fclose(fp);
172 }
173 }
174 if (!cacheLoaded) {
175 if (GetProcessingSettings().debugLevel >= 0) {
176 GPUInfo("Starting CUDA RTC Compilation");
177 }
178 HighResTimer rtcTimer;
179 rtcTimer.ResetStart();
180 tbb::parallel_for<uint32_t>(0, nCompile, [&](auto i) {
181 if (GetProcessingSettings().debugLevel >= 3) {
182 printf("Compiling %s\n", (filename + "_" + std::to_string(i) + mRtcSrcExtension).c_str());
183 }
184 FILE* fp = fopen((filename + "_" + std::to_string(i) + mRtcSrcExtension).c_str(), "w+b");
185 if (fp == nullptr) {
186 throw std::runtime_error("Error opening file");
187 }
188
189 std::string kernel = "extern \"C\" {";
190 kernel += GetProcessingSettings().rtc.compilePerKernel ? kernels[i] : kernelsall;
191 kernel += "}";
192
193 bool deterministic = GetProcessingSettings().rtc.deterministic || (GetProcessingSettings().rtc.compilePerKernel && o2::gpu::internal::noFastMathKernels.find(GetKernelName(i)) != o2::gpu::internal::noFastMathKernels.end());
194 const std::string deterministicStr = std::string(deterministic ? "#define GPUCA_DETERMINISTIC_CODE(det, indet) det\n" : "#define GPUCA_DETERMINISTIC_CODE(det, indet) indet\n");
195
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");
202 }
203 fclose(fp);
204 std::string command = baseCommand;
205 if (deterministic) {
206 command += std::string(" ") + std::string(_binary_GPUReconstructionCUDArtc_command_no_fast_math_start, _binary_GPUReconstructionCUDArtc_command_no_fast_math_len);
207 }
208 command += " -c " + filename + "_" + std::to_string(i) + mRtcSrcExtension + " -o " + filename + "_" + std::to_string(i) + mRtcBinExtension;
209 if (GetProcessingSettings().debugLevel < 0) {
210 command += " &> /dev/null";
211 } else if (GetProcessingSettings().debugLevel < 2) {
212 command += " > /dev/null";
213 }
214 if (GetProcessingSettings().debugLevel >= 3) {
215 printf("Running command %s\n", command.c_str());
216 }
217 if (system(command.c_str())) {
218 if (GetProcessingSettings().debugLevel >= 3) {
219 printf("Source code file: %s", filename.c_str());
220 }
221 throw std::runtime_error("Error during CUDA compilation");
222 } // clang-format off
223 }, tbb::simple_partitioner()); // clang-format on
224 if (GetProcessingSettings().debugLevel >= 0) {
225 GPUInfo("RTC Compilation finished (%f seconds)", rtcTimer.GetCurrentElapsedTime());
226 }
227 if (GetProcessingSettings().rtc.cacheOutput) {
228 FILE* fp = fopen((GetProcessingSettings().rtctech.cacheFolder + "/rtc.cuda.cache").c_str(), "w+b");
229 if (fp == nullptr) {
230 throw std::runtime_error("Cannot open cache file for writing");
231 }
232 GPUInfo("Storing RTC compilation result in cache file");
233
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 ||
239 fwrite(&GetProcessingSettings().rtc, sizeof(GetProcessingSettings().rtc), 1, fp) != 1) {
240 throw std::runtime_error("Error writing cache file");
241 }
242
243 std::vector<char> buffer;
244 for (uint32_t i = 0; i < nCompile; i++) {
245 FILE* fp2 = fopen((filename + "_" + std::to_string(i) + mRtcBinExtension).c_str(), "rb");
246 if (fp2 == nullptr) {
247 throw std::runtime_error("Cannot open cuda module file");
248 }
249 fseek(fp2, 0, SEEK_END);
250 size_t size = ftell(fp2);
251 buffer.resize(size);
252 fseek(fp2, 0, SEEK_SET);
253 if (fread(buffer.data(), 1, size, fp2) != size) {
254 throw std::runtime_error("Error reading cuda module file");
255 }
256 fclose(fp2);
257
258 if (fwrite(&size, sizeof(size), 1, fp) != 1 ||
259 fwrite(buffer.data(), 1, size, fp) != size) {
260 throw std::runtime_error("Error writing cache file");
261 }
262 }
263 fclose(fp);
264 }
265 }
266 if (GetProcessingSettings().rtc.cacheOutput && GetProcessingSettings().rtctech.cacheMutex) {
267 if (lockf(fd, F_ULOCK, 0)) {
268 throw std::runtime_error("Error unlocking RTC cache mutex file");
269 }
270 close(fd);
271 }
272
273 return 0;
274}
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)