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