1 /*
2 * Copyright (c) 2024, 2026, Oracle and/or its affiliates. All rights reserved.
3 * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
4 *
5 * This code is free software; you can redistribute it and/or modify it
6 * under the terms of the GNU General Public License version 2 only, as
7 * published by the Free Software Foundation. Oracle designates this
8 * particular file as subject to the "Classpath" exception as provided
9 * by Oracle in the LICENSE file that accompanied this code.
10 *
11 * This code is distributed in the hope that it will be useful, but WITHOUT
12 * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
13 * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
14 * version 2 for more details (a copy is included in the LICENSE file that
15 * accompanied this code).
16 *
17 * You should have received a copy of the GNU General Public License version
18 * 2 along with this work; if not, write to the Free Software Foundation,
19 * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA.
20 *
21 * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA
22 * or visit www.oracle.com if you need additional information or have any
23 * questions.
24 */
25
26 #include <sys/wait.h>
27 #include <unistd.h>
28 #include <chrono>
29 #include "cuda_backend.h"
30 #include <iostream>
31 #include <cstdlib>
32 #include <cstring>
33 #include <dlfcn.h>
34 #include <filesystem>
35 #include <sstream>
36 #include <string>
37 #include <vector>
38
39 namespace {
40
41 bool isCudaTileIR(const char *image, size_t len) {
42 // The Tile IR bytecode magic number consumes 8 bytes and is \x7FTileIR\x00.
43 // The magic number must be present at the beginning of the bytecode file to
44 // be accepted by the driver.
45 // See: https://docs.nvidia.com/cuda/tile-ir/13.3/sections/bytecode.html#magic-number
46 static constexpr char kMagic[] = "\x7f" "TileIR";
47 return image != nullptr && len >= sizeof(kMagic) &&
48 std::memcmp(image, kMagic, sizeof(kMagic)) == 0;
49 }
50
51 // Tile code generation starts at sm_80 for both nvcc and NVRTC. sm_7x,
52 // including the nvcc default sm_75, does not generate Tile code.
53 // https://docs.nvidia.com/cuda/cuda-tile-cpp-api-reference/general_principles.html
54 void assertTileArchSupport(int major, int minor, const std::string &cudaPath) {
55 if (major >= 8) {
56 return;
57 }
58 std::cerr << "Tile compilation requires sm_80 or newer, got sm_"
59 << major << minor << ". CUDA source saved to " << cudaPath
60 << std::endl;
61 std::exit(1);
62 }
63
64 bool isLibjsigLoaded() {
65 static const bool loaded = [] {
66 void *handle = dlopen(nullptr, RTLD_LAZY);
67 if (handle == nullptr) {
68 return false;
69 }
70 dlerror();
71 const bool found =
72 dlsym(handle, "JVM_begin_signal_setting") != nullptr;
73 dlclose(handle);
74 return found;
75 }();
76 return loaded;
77 }
78
79 void warnNvrtcRequiresLibjsig() {
80 static bool warned = false;
81 if (warned) {
82 return;
83 }
84 warned = true;
85 std::cout << "[HAT] WARNING: libjsig not preloaded. Loading Tile IR "
86 << "in-process processes signals first and resends them with a "
87 << "polluted state, breaking JVM signal handling. NVRTC emits "
88 << "only Tile IR for Tile kernels, so libjsig is required. "
89 << "Falling back to nvcc for Tile. This warning is issued once."
90 << "To use NVRTC Tile, launch with:"
91 << std::endl;
92 if (const char *javaHome = std::getenv("JAVA_HOME");
93 javaHome != nullptr && javaHome[0] != '\0') {
94 std::cout << " LD_PRELOAD=" << javaHome << "/lib/libjsig.so java ..."
95 << std::endl;
96 } else {
97 std::cout << " LD_PRELOAD=$JAVA_HOME/lib/libjsig.so java ..."
98 << std::endl;
99 }
100 }
101
102 using nvrtcProgram = void *;
103 using nvrtcResult = int;
104 constexpr nvrtcResult NVRTC_SUCCESS = 0;
105
106 template <typename T>
107 T loadNvrtcSymbol(void *handle, const char *name, bool required = true) {
108 dlerror();
109 void *symbol = dlsym(handle, name);
110 if (const char *error = dlerror()) {
111 if (!required) {
112 return nullptr;
113 }
114 std::cerr << "Failed to load NVRTC symbol " << name << ": " << error << std::endl;
115 std::exit(1);
116 }
117 return reinterpret_cast<T>(symbol);
118 }
119
120 void *loadNvrtcLibrary() {
121 std::vector<std::string> candidates;
122 if (const char *envLibrary = std::getenv("HAT_CUDA_NVRTC_LIBRARY")) {
123 dlerror();
124 if (void *handle = dlopen(envLibrary, RTLD_NOW | RTLD_LOCAL)) {
125 return handle;
126 }
127 std::cerr << "Failed to load NVRTC from HAT_CUDA_NVRTC_LIBRARY='"
128 << envLibrary << "'";
129 if (const char *error = dlerror()) {
130 std::cerr << ": " << error;
131 }
132 std::cerr << std::endl;
133 std::exit(1);
134 }
135 // First try the CUDA toolkit library directory found by CMake.
136 // Then let dlopen search the system loader path.
137 #ifdef HAT_CUDA_LIBRARY_DIR
138 candidates.emplace_back(std::string(HAT_CUDA_LIBRARY_DIR) +
139 "/libnvrtc.so");
140 #endif
141 candidates.emplace_back("libnvrtc.so");
142
143 std::vector<std::string> errors;
144 for (const std::string &candidate : candidates) {
145 if (void *handle = dlopen(candidate.c_str(), RTLD_NOW | RTLD_LOCAL)) {
146 return handle;
147 }
148 if (const char *error = dlerror()) {
149 errors.emplace_back(candidate + ": " + error);
150 }
151 }
152
153 std::cerr << "Failed to load NVRTC. Set HAT_CUDA_NVRTC_LIBRARY to "
154 << "the NVRTC shared library path or name, "
155 << "or use HAT=CUDA_NVCC." << std::endl;
156 for (const std::string &error : errors) {
157 std::cerr << " " << error << std::endl;
158 }
159 std::exit(1);
160 }
161
162 std::vector<std::string> splitDelimited(const char *value,
163 const char delimiter) {
164 std::vector<std::string> result;
165 if (value == nullptr || *value == '\0') {
166 return result;
167 }
168
169 std::stringstream stream(value);
170 std::string item;
171 while (std::getline(stream, item, delimiter)) {
172 if (!item.empty()) {
173 result.push_back(item);
174 }
175 }
176 return result;
177 }
178
179 // Minimal NVRTC ABI surface used through dlsym. This avoids a build-time
180 // dependency on nvrtc.h and a link-time dependency on libnvrtc.so.
181 struct NvrtcApi {
182 using CreateProgram = nvrtcResult (*)(nvrtcProgram *,
183 const char *,
184 const char *,
185 int,
186 const char * const *,
187 const char * const *);
188 using CompileProgram = nvrtcResult (*)(nvrtcProgram, int, const char * const *);
189 using DestroyProgram = nvrtcResult (*)(nvrtcProgram *);
190 using GetErrorString = const char *(*)(nvrtcResult);
191 using GetPTX = nvrtcResult (*)(nvrtcProgram, char *);
192 using GetPTXSize = nvrtcResult (*)(nvrtcProgram, size_t *);
193 using GetProgramLog = nvrtcResult (*)(nvrtcProgram, char *);
194 using GetProgramLogSize = nvrtcResult (*)(nvrtcProgram, size_t *);
195 using GetTileIR = nvrtcResult (*)(nvrtcProgram, char *);
196 using GetTileIRSize = nvrtcResult (*)(nvrtcProgram, size_t *);
197
198 void *handle;
199 CreateProgram createProgram;
200 CompileProgram compileProgram;
201 DestroyProgram destroyProgram;
202 GetErrorString getErrorString;
203 GetPTX getPTX;
204 GetPTXSize getPTXSize;
205 GetProgramLog getProgramLog;
206 GetProgramLogSize getProgramLogSize;
207 GetTileIR getTileIR;
208 GetTileIRSize getTileIRSize;
209
210 explicit NvrtcApi(void *handle)
211 : handle(handle),
212 createProgram(loadNvrtcSymbol<CreateProgram>(handle, "nvrtcCreateProgram")),
213 compileProgram(loadNvrtcSymbol<CompileProgram>(handle, "nvrtcCompileProgram")),
214 destroyProgram(loadNvrtcSymbol<DestroyProgram>(handle, "nvrtcDestroyProgram")),
215 getErrorString(loadNvrtcSymbol<GetErrorString>(handle, "nvrtcGetErrorString")),
216 getPTX(loadNvrtcSymbol<GetPTX>(handle, "nvrtcGetPTX")),
217 getPTXSize(loadNvrtcSymbol<GetPTXSize>(handle, "nvrtcGetPTXSize")),
218 getProgramLog(loadNvrtcSymbol<GetProgramLog>(handle, "nvrtcGetProgramLog")),
219 getProgramLogSize(loadNvrtcSymbol<GetProgramLogSize>(handle, "nvrtcGetProgramLogSize")),
220 getTileIR(loadNvrtcSymbol<GetTileIR>(handle, "nvrtcGetTileIR", false)),
221 getTileIRSize(loadNvrtcSymbol<GetTileIRSize>(handle, "nvrtcGetTileIRSize", false)) {
222 }
223 };
224
225 // Load and cache NVRTC API entry points on first use
226 NvrtcApi &nvrtcApi() {
227 static NvrtcApi api(loadNvrtcLibrary());
228 return api;
229 }
230
231 void nvrtcCheck(const NvrtcApi &api, const nvrtcResult result,
232 const char *functionName) {
233 if (result != NVRTC_SUCCESS) {
234 std::cerr << functionName << " NVRTC error = " << result << " "
235 << api.getErrorString(result) << std::endl;
236 std::exit(1);
237 }
238 }
239
240 std::string getNvrtcLog(const NvrtcApi &api, nvrtcProgram program) {
241 size_t logSize = 0;
242 nvrtcCheck(api,
243 api.getProgramLogSize(program, &logSize),
244 "nvrtcGetProgramLogSize");
245 std::string log;
246 if (logSize > 1) {
247 log.resize(logSize, '\0');
248 nvrtcCheck(api,
249 api.getProgramLog(program, log.data()),
250 "nvrtcGetProgramLog");
251 }
252 return log;
253 }
254 }
255
256 CudaImage::CudaImage()
257 : Text(0L) {
258 }
259
260 CudaImage::CudaImage(size_t len)
261 : Text(len) {
262 }
263
264 CudaImage::CudaImage(char *text)
265 : Text(text, false) {
266 }
267
268 CudaImage::CudaImage(size_t len, char *text)
269 : Text(len, text, true) {
270 }
271 CudaImage::CudaImage(size_t len, char *text, bool isCopy)
272 : Text(len, text, isCopy) {
273 }
274
275 CudaSource::CudaSource(size_t len)
276 : Text(len) {
277 }
278
279 CudaSource::CudaSource(char *text)
280 : Text(text, false) {
281 }
282
283 CudaSource::CudaSource(size_t len, char *text, bool isCopy, bool lineinfo, int typeModel)
284 : Text(len, text, isCopy) {
285 _lineInfo = lineinfo;
286 _typeModel = typeModel;
287 }
288
289 CudaSource::CudaSource()
290 : Text(0) {
291 }
292
293 bool CudaSource::lineInfo() const {
294 return _lineInfo;
295 }
296
297 int CudaSource::typeModel() const {
298 return _typeModel;
299 }
300
301 uint64_t timeSinceEpochMillisec() {
302 using namespace std::chrono;
303 return duration_cast<milliseconds>(system_clock::now().time_since_epoch()).count();
304 }
305
306 std::string tmpFileName(uint64_t time, const std::string directoryName, const std::string &suffix) {
307 std::stringstream timestamp;
308 timestamp << directoryName << "/tmp_" << time << suffix;
309 return timestamp.str();
310 }
311
312 CudaBackend::CudaBackend(int configBits)
313 : Backend(new Config(configBits), new CudaQueue(this)),
314 initStatus(cuInit(0)), device(), context(), major(0), minor(0) {
315 int deviceCount = 0;
316
317 if (initStatus == CUDA_SUCCESS) {
318 CUDA_CHECK(cuDeviceGetCount(&deviceCount), "cuDeviceGetCount");
319 if (config->info) {
320 std::cout << "CudaBackend device count = " << deviceCount << std::endl;
321 }
322 CUDA_CHECK(cuDeviceGet(&device, 0), "cuDeviceGet");
323 CUDA_CHECK(cuDeviceGetAttribute(&major,
324 CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR,
325 device),
326 "cuDeviceGetAttribute");
327 CUDA_CHECK(cuDeviceGetAttribute(&minor,
328 CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR,
329 device),
330 "cuDeviceGetAttribute");
331 #if defined(CUDA_VERSION) && CUDA_VERSION >= 12080
332 CUctxCreateParams ctxCreateParams = {};
333 CUDA_CHECK(cuCtxCreate_v4(&context, &ctxCreateParams, 0, device), "cuCtxCreate");
334 #else
335 // Invoke previous implementation with 3 parameters
336 CUDA_CHECK(cuCtxCreate(&context, 0, device), "cuCtxCreate");
337 #endif
338 if (config->info) {
339 std::cout << "CudaBackend context created ok (id=" << context << ")" << std::endl;
340 }
341 dynamic_cast<CudaQueue *>(queue)->init();
342 } else {
343 CUDA_CHECK(initStatus, "cuInit() failed we seem to have the runtime library but no device");
344 }
345 }
346
347 CudaBackend::~CudaBackend() {
348 std::cout << "freeing context" << std::endl;
349 CUDA_CHECK(cuCtxDestroy(context), "cuCtxDestroy");
350 }
351
352 void CudaBackend::shortDeviceInfo() {
353 char name[100];
354 CUDA_CHECK(cuDeviceGetName(name, sizeof(name), device), "cuDeviceGetName");
355 std::cout << "[INFO] Using NVIDIA GPU: " << name << std::endl;
356 }
357
358 void CudaBackend::showDeviceInfo() {
359 char name[100];
360 CUDA_CHECK(cuDeviceGetName(name, sizeof(name), device), "cuDeviceGetName");
361
362 std::cout << "> Using device 0: " << name << std::endl;
363 std::cout << "> GPU Device has major=" << major << " minor=" << minor << " compute capability" << std::endl;
364
365 int warpSize;
366 CUDA_CHECK(cuDeviceGetAttribute(&warpSize, CU_DEVICE_ATTRIBUTE_WARP_SIZE, device), "cuDeviceGetAttribute");
367 std::cout << "> GPU Device has warpSize " << warpSize << std::endl;
368
369 int threadsPerBlock;
370 CUDA_CHECK(cuDeviceGetAttribute(&threadsPerBlock, CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK, device), "cuDeviceGetAttribute");
371 std::cout << "> GPU Device has threadsPerBlock " << threadsPerBlock << std::endl;
372
373 int cores;
374 CUDA_CHECK(cuDeviceGetAttribute(&cores, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, device), "cuDeviceGetAttribute");
375 std::cout << "> GPU Cores " << cores << std::endl;
376
377 size_t totalGlobalMem;
378 CUDA_CHECK(cuDeviceTotalMem(&totalGlobalMem, device), "cuDeviceTotalMem");
379 std::cout << " Total amount of global memory: " << (unsigned long long) totalGlobalMem << std::endl;
380 std::cout << " 64-bit Memory Address: " <<
381 ((totalGlobalMem > static_cast<unsigned long long>(4) * 1024 * 1024 * 1024L) ? "YES" : "NO") << std::endl;
382 }
383
384 // TypeModel = 0; SIMT model
385 // TypeModel = 1; Tile model
386 bool CudaBackend::useNvrtcCompiler(int typeModel) const {
387 // CUDA NVRTC is the default. Thus, to obtain nvrtc we need
388 // to check that NVCC flag is not enabled.
389 const bool useNVCC = config->cudaNvcc;
390
391 // Loading Tile IR in-process registers handlers that process signals first
392 // and resend them with a polluted state, breaking JVM signal handling.
393 // NVRTC emits only Tile IR for Tile kernels, so libjsig is required.
394 // Tile cubin and SIMT PTX do not need it.
395 if (!useNVCC && typeModel > 0 && !isLibjsigLoaded()) {
396 warnNvrtcRequiresLibjsig();
397 return false;
398 }
399 // use NVRTC is the flag NVCC is not enabled.
400 return !useNVCC;
401 }
402
403 CudaImage *CudaBackend::nvcc(const CudaSource *cudaSource) {
404 // create var/cuda directory
405 std::string localDirectory = "./var/cuda";
406 std::filesystem::create_directories(localDirectory);
407 // create temp file for cuda generated code
408 const uint64_t time = timeSinceEpochMillisec();
409 std::string suffix = ".ptx";
410 if (cudaSource->typeModel() > 0) {
411 suffix = ".cubin";
412 }
413 const std::string ptxPath = tmpFileName(time, localDirectory, suffix);
414 const std::string cudaPath = tmpFileName(time, localDirectory, ".cu");
415 cudaSource->write(cudaPath);
416
417 if (cudaSource->typeModel() > 0) {
418 assertTileArchSupport(major, minor, cudaPath);
419 }
420 const std::string smVersion =
421 std::string("sm_") + std::to_string(major) + std::to_string(minor);
422
423 // compile the generated code
424 int pid;
425 if ((pid = fork()) == 0) { //child
426 const auto cudaCompiler = "nvcc";
427 std::vector<std::string> command;
428 command.push_back(cudaCompiler);
429
430 if (cudaSource->typeModel() > 0) {
431 command.push_back("--tilecubin");
432 command.push_back("--std=c++20");
433 command.push_back("--enable-tile");
434 } else {
435 command.push_back("-ptx");
436 }
437 command.push_back("-arch");
438 command.push_back(smVersion);
439
440 command.push_back(cudaPath);
441 if (cudaSource->lineInfo()) {
442 command.push_back("-lineinfo");
443 }
444
445 command.push_back("-o");
446 command.push_back(ptxPath);
447
448 // conver to char*[] fr the execvp function
449 const char* args[command.size() + 1];
450 for (int i = 0; i < command.size(); i++) {
451 args[i] = command[i].c_str();
452 }
453 args[command.size()] = nullptr;
454 const int stat = execvp(cudaCompiler, (char *const *) args);
455 std::cerr << " nvcc stat = " << stat << " errno=" << errno << " '" << std::strerror(errno) << "'" << std::endl;
456 std::exit(errno);
457 } else if (pid < 0) {// fork failed.
458 std::cerr << "fork of nvcc failed" << std::endl;
459 std::exit(1);
460 } else { //parent
461 int status;
462 pid_t result = wait(&status);
463 auto *ptx = new CudaImage();
464 ptx->read(ptxPath);
465 return ptx;
466 }
467 }
468
469 CudaImage *CudaBackend::nvrtc(const CudaSource *cudaSource) {
470 NvrtcApi &api = nvrtcApi();
471 const bool tile = cudaSource->typeModel() > 0;
472 std::string source(cudaSource->text, cudaSource->len);
473
474 // Keep generated CUDA/PTX artifacts under the same directory as the nvcc path.
475 std::string localDirectory = "./var/cuda";
476 std::filesystem::create_directories(localDirectory);
477 const uint64_t time = timeSinceEpochMillisec();
478 const std::string imagePath =
479 tmpFileName(time, localDirectory, tile ? ".tileir" : ".ptx");
480 const std::string cudaPath = tmpFileName(time, localDirectory, ".cu");
481 cudaSource->write(cudaPath);
482
483 if (tile && (api.getTileIR == nullptr || api.getTileIRSize == nullptr)) {
484 std::cerr << "NVRTC Tile compilation requires CUDA Toolkit 13.3 or later "
485 << "(nvrtcGetTileIR is missing). Use HAT=CUDA_NVCC, "
486 << "or upgrade the NVRTC library." << std::endl;
487 std::exit(1);
488 }
489
490 if (tile) {
491 assertTileArchSupport(major, minor, cudaPath);
492 }
493
494 nvrtcProgram program;
495 nvrtcCheck(api,
496 api.createProgram(&program,
497 source.c_str(),
498 cudaPath.c_str(),
499 0,
500 nullptr,
501 nullptr),
502 "nvrtcCreateProgram");
503
504 std::vector<std::string> options;
505 if (tile) {
506 options.emplace_back("--std=c++20");
507 options.emplace_back("--enable-tile");
508 } else {
509 options.emplace_back("--std=c++17");
510 }
511 options.emplace_back("--gpu-architecture=compute_" +
512 std::to_string(major) +
513 std::to_string(minor));
514 #ifdef HAT_CUDA_INCLUDE_DIRS
515 // CMake serializes the CUDA include list with '|' so it can be passed as a
516 // single compile definition.
517 for (const std::string &includeDir :
518 splitDelimited(HAT_CUDA_INCLUDE_DIRS, '|')) {
519 options.emplace_back("-I" + includeDir);
520 }
521 #endif
522 if (cudaSource->lineInfo()) {
523 options.emplace_back("--generate-line-info");
524 }
525
526 std::vector<const char *> optionPtrs;
527 optionPtrs.reserve(options.size());
528 for (const std::string &option : options) {
529 optionPtrs.push_back(option.c_str());
530 }
531
532 nvrtcResult compileResult =
533 api.compileProgram(program,
534 static_cast<int>(optionPtrs.size()),
535 optionPtrs.data());
536 if (compileResult != NVRTC_SUCCESS) {
537 std::cerr << "NVRTC compilation failed: "
538 << api.getErrorString(compileResult)
539 << ". CUDA source saved to " << cudaPath << std::endl;
540 std::string log = getNvrtcLog(api, program);
541 if (!log.empty()) {
542 std::cerr << "> NVRTC log:" << std::endl << log << std::endl;
543 }
544 nvrtcCheck(api, api.destroyProgram(&program), "nvrtcDestroyProgram");
545 std::exit(1);
546 } else if (config->info || config->trace) {
547 std::string log = getNvrtcLog(api, program);
548 if (!log.empty()) {
549 std::cout << "> NVRTC log:" << std::endl << log << std::endl;
550 }
551 }
552
553 size_t imageSize = 0;
554 if (tile) {
555 nvrtcCheck(api, api.getTileIRSize(program, &imageSize), "nvrtcGetTileIRSize");
556 if (imageSize == 0) {
557 std::cerr << "NVRTC returned empty Tile IR. CUDA source saved to "
558 << cudaPath << std::endl;
559 nvrtcCheck(api, api.destroyProgram(&program), "nvrtcDestroyProgram");
560 std::exit(1);
561 }
562 auto *image = new CudaImage(imageSize);
563 nvrtcCheck(api, api.getTileIR(program, image->text), "nvrtcGetTileIR");
564 image->write(imagePath);
565 nvrtcCheck(api, api.destroyProgram(&program), "nvrtcDestroyProgram");
566 return image;
567 }
568
569 nvrtcCheck(api, api.getPTXSize(program, &imageSize), "nvrtcGetPTXSize");
570 auto *ptx = new CudaImage(imageSize);
571 nvrtcCheck(api, api.getPTX(program, ptx->text), "nvrtcGetPTX");
572 if (imageSize == 0 || ptx->text[imageSize - 1] != '\0') {
573 std::cerr << "NVRTC returned invalid PTX buffer" << std::endl;
574 nvrtcCheck(api, api.destroyProgram(&program), "nvrtcDestroyProgram");
575 std::exit(1);
576 }
577
578 // nvrtcGetPTXSize includes the trailing NUL. Keep it in memory for the
579 // driver API, but omit it from the debug artifact to match nvcc output.
580 CudaImage ptxFile(imageSize > 0 ? imageSize - 1 : 0, ptx->text, false);
581 ptxFile.write(imagePath);
582
583 nvrtcCheck(api, api.destroyProgram(&program), "nvrtcDestroyProgram");
584 return ptx;
585 }
586
587 CudaBackend::CudaModule *CudaBackend::compile(const CudaSource &cudaSource) {
588 return compile(&cudaSource);
589 }
590
591 CudaBackend::CudaModule *CudaBackend::compile(const CudaSource *cudaSource) {
592 const bool useNvrtc = useNvrtcCompiler(cudaSource->typeModel());
593 if (config->info) {
594 std::cout << "[INFO] CUDA source compiler: "
595 << (useNvrtc ? "NVRTC" : "NVCC") << std::endl;
596 }
597 const CudaImage *image = useNvrtc ? nvrtc(cudaSource) : nvcc(cudaSource);
598 return compile(image);
599 }
600
601 CudaBackend::CudaModule *CudaBackend::compile(const CudaImage &image) {
602 return compile(&image);
603 }
604
605 CudaBackend::CudaModule *CudaBackend::compile(const CudaImage *image) {
606 CUmodule module;
607 if (image->text != nullptr) {
608 CUDA_CHECK(cuCtxSetCurrent(context), "cuCtxSetCurrent");
609 if (isCudaTileIR(image->text, image->len)) {
610 CUDA_CHECK(cuModuleLoadData(&module, image->text), "cuModuleLoadData");
611 return new CudaModule(this, image, new char[1]{'\0'}, true, module);
612 }
613
614 const Log *infLog = new Log(8192);
615 const Log *errLog = new Log(8192);
616 constexpr unsigned int optc = 5;
617 const auto jitOptions = new CUjit_option[optc];
618 auto jitOptVals = new void *[optc];
619
620 jitOptions[0] = CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES;
621 jitOptVals[0] = reinterpret_cast<void *>(infLog->len);
622 jitOptions[1] = CU_JIT_INFO_LOG_BUFFER;
623 jitOptVals[1] = infLog->text;
624 jitOptions[2] = CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES;
625 jitOptVals[2] = reinterpret_cast<void *>(errLog->len);
626 jitOptions[3] = CU_JIT_ERROR_LOG_BUFFER;
627 jitOptVals[3] = errLog->text;
628 jitOptions[4] = CU_JIT_GENERATE_LINE_INFO;
629 jitOptVals[4] = reinterpret_cast<void *>(1);
630
631 CUDA_CHECK(cuModuleLoadDataEx(&module, image->text, optc, jitOptions, (void **) jitOptVals), "cuModuleLoadDataEx");
632
633 if (*infLog->text!='\0'){
634 std::cout << "> PTX JIT inflog:" << std::endl << infLog->text << std::endl;
635 }
636 if (*errLog->text!='\0'){
637 std::cout << "> PTX JIT errlog:" << std::endl << errLog->text << std::endl;
638 }
639 return new CudaModule(this, image, infLog->text, true, module);
640 //delete ptx;
641 }
642 std::cout << "no ptx content!" << std::endl;
643 exit(1);
644 }
645
646 //Entry point from HAT. We use the config PTX bit to determine which Source type
647 Backend::CompilationUnit *CudaBackend::compile(const int len, char *source, int typeModel) {
648 if (config->traceCalls) {
649 std::cout << "inside compileProgram" << std::endl;
650 }
651
652 if (config->ptx){
653 if (config->trace) {
654 std::cout << "compiling from provided ptx " << std::endl;
655 }
656 CudaImage image(len, source, false);
657 return compile(image);
658 }else{
659 if (config->trace) {
660 std::cout << "compiling from provided cuda " << std::endl;
661 }
662 CudaSource cudaSource(len , source, false, config->profile, typeModel);
663 return compile(cudaSource);
664 }
665 }
666
667 /*
668
669 if (config->ptx) {
670
671 } else {
672 if (config->trace) {
673 std::cout << "compiling from cuda c99 " << std::endl;
674 }
675 if (config->showCode) {
676 std::cout << "cuda " << source << std::endl;
677 }
678 auto* cuda = new CudaSource(len, source, false);
679 ptx = nvcc(cuda);
680 }
681 if (config->showCode) {
682 std::cout << "ptx " << ptx->text << std::endl;
683 }
684 CUmodule module;
685
686
687 if (ptx->text != nullptr) {
688 constexpr unsigned int jitNumOptions = 2;
689 const auto jitOptions = new CUjit_option[jitNumOptions];
690 const auto jitOptVals = new void *[jitNumOptions];
691
692 // set up size of compilation log buffer
693 jitOptions[0] = CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES;
694 constexpr int jitLogBufferSize = 8192;
695 jitOptVals[0] = reinterpret_cast<void *>(jitLogBufferSize);
696
697 // set up pointer to the compilation log buffer
698 jitOptions[1] = CU_JIT_INFO_LOG_BUFFER;
699 auto jitLogBuffer = new char[jitLogBufferSize];
700 jitOptVals[1] = jitLogBuffer;
701 cuCtxSetCurrent(context);
702
703 WHERE{
704 .f = __FILE__, .l = __LINE__,
705 .e = cuModuleLoadDataEx(&module, ptx->text, jitNumOptions, jitOptions, jitOptVals),
706 .t = "cuModuleLoadDataEx"
707 }.report();
708 if (jitLogBuffer != nullptr && *jitLogBuffer!='\0'){
709 std::cout << "PTX log:" << jitLogBuffer << std::endl;
710 }
711 return new CudaModule(this, ptx->text, jitLogBuffer, true, module);
712 } else {
713 std::cout << "no ptx content!" << std::endl;
714 exit(1);
715 }
716 } */
717
718 extern "C" long getBackend(int mode) {
719 long backendHandle = reinterpret_cast<long>(new CudaBackend(mode));
720 // std::cout << "getBackend() -> backendHandle=" << std::hex << backendHandle << std::dec << std::endl;
721 return backendHandle;
722 }
723
724 void clCallback(void *) {
725 std::cerr << "start of compute" << std::endl;
726 }
727
728 void CudaBackend::computeEnd() {
729 queue->computeEnd();
730 }
731
732 void CudaBackend::computeStart() {
733 queue->computeStart();
734 }
735
736 bool CudaBackend::getBufferFromDeviceIfDirty(void *memorySegment, long memorySegmentLength) {
737 if (config->traceCalls) {
738 std::cout << "getBufferFromDeviceIfDirty(" << std::hex << reinterpret_cast<long>(memorySegment) << "," <<
739 std::dec << memorySegmentLength << "){" << std::endl;
740 }
741 if (config->minimizeCopies) {
742 const BufferState *bufferState = BufferState::of(memorySegment, memorySegmentLength);
743 if (bufferState->state == BufferState::DEVICE_OWNED) {
744 queue->copyFromDevice(static_cast<Backend::Buffer *>(bufferState->vendorPtr));
745 if (config->traceEnqueues | config->traceCopies) {
746 std::cout << "copying buffer from device (from java access) " << std::endl;
747 }
748 queue->wait();
749 queue->release();
750 } else {
751 std::cout << "HOW DID WE GET HERE 1 attempting to get buffer but buffer is not device dirty" << std::endl;
752 std::exit(1);
753 }
754 } else {
755 std::cerr <<
756 "HOW DID WE GET HERE ? java side should avoid calling getBufferFromDeviceIfDirty as we are not minimising buffers!"
757 << std::endl;
758 std::exit(1);
759 }
760 if (config->traceCalls) {
761 std::cout << "}getBufferFromDeviceIfDirty()" << std::endl;
762 }
763 return true;
764 }
765
766 CudaBackend *CudaBackend::of(const long backendHandle) {
767 return reinterpret_cast<CudaBackend *>(backendHandle);
768 }
769
770 CudaBackend *CudaBackend::of(Backend *backend) {
771 return dynamic_cast<CudaBackend *>(backend);
772 }
773
774 CudaBackend::CudaBuffer *CudaBackend::getOrCreateBuffer(BufferState *bufferState) {
775 CudaBuffer *cudaBuffer = nullptr;
776 if (bufferState->vendorPtr == nullptr || bufferState->state == BufferState::NEW_STATE) {
777 cudaBuffer = new CudaBuffer(this, bufferState);
778 if (config->trace) {
779 std::cout << "We allocated arg buffer " << std::endl;
780 }
781 bufferState->state = BufferState::NEW_STATE;
782 } else {
783 if (config->trace) {
784 std::cout << "Were reusing buffer buffer " << std::endl;
785 }
786 cudaBuffer = static_cast<CudaBuffer *>(bufferState->vendorPtr);
787 }
788 return cudaBuffer;
789 }