Lines Matching +full:- +full:s
1 // Copyright (c) 2017-2026, Lawrence Livermore National Security, LLC and other CEED contributors.
2 // All Rights Reserved. See the top-level LICENSE and NOTICE files for details.
4 // SPDX-License-Identifier: BSD-2-Clause
8 #include "ceed-cuda-compile.h"
12 #include <ceed/jit-tools.h>
28 #include "ceed-cuda-common.h"
44 //------------------------------------------------------------------------------
46 //------------------------------------------------------------------------------
48 CeedDebug(ceed, "Running command:\n$ %s", command); in CeedCallSystem_Core()
51 …CeedCheck(output_stream != nullptr, ceed, CEED_ERROR_BACKEND, "Failed to %s\ncommand:\n$ %s", mess… in CeedCallSystem_Core()
59 CeedDebug(ceed, "output:\n%s\n", output.c_str()); in CeedCallSystem_Core()
60 …se(output_stream) == 0, ceed, CEED_ERROR_BACKEND, "Failed to %s\ncommand:\n$ %s\nerror:\n%s", mess… in CeedCallSystem_Core()
64 //------------------------------------------------------------------------------
66 //------------------------------------------------------------------------------
107 code << "#include <ceed/jit-source/cuda/cuda-jit.h>\n\n"; in CeedCompileCore_Cuda()
109 // Non-macro options in CeedCompileCore_Cuda()
111 opts[0] = "-default-device"; in CeedCompileCore_Cuda()
113 CeedCallCuda(ceed, cudaGetDeviceProperties(&prop, ceed_data->device_id)); in CeedCompileCore_Cuda()
117 // -arch, since it was only emitting PTX. It will now support actual in CeedCompileCore_Cuda()
119 // https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#dynamic-code-generation in CeedCompileCore_Cuda()
120 "-arch=sm_" in CeedCompileCore_Cuda()
122 "-arch=compute_" in CeedCompileCore_Cuda()
126 opts[2] = "-Dint32_t=int"; in CeedCompileCore_Cuda()
127 opts[3] = "-DCEED_RUNNING_JIT_PASS=1"; in CeedCompileCore_Cuda()
137 include_dir_arg << "-I" << jit_source_dirs[i]; in CeedCompileCore_Cuda()
151 define_arg << "-D" << jit_defines[i]; in CeedCompileCore_Cuda()
161 …CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "---------- ATTEMPTING TO COMPILE JIT SOURCE --------… in CeedCompileCore_Cuda()
162 CeedDebug(ceed, "Source:\n%s\n", code.str().c_str()); in CeedCompileCore_Cuda()
163 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "---------- END OF JIT SOURCE ----------\n"); in CeedCompileCore_Cuda()
170 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "---------- JiT COMPILER OPTIONS ----------\n"); in CeedCompileCore_Cuda()
172 CeedDebug(ceed, "Option %d: %s", i, opts[i]); in CeedCompileCore_Cuda()
175 …CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "---------- END OF JiT COMPILER OPTIONS ----------\n"… in CeedCompileCore_Cuda()
197 return CeedError(ceed, CEED_ERROR_BACKEND, "%s\n%s", nvrtcGetErrorString(result), log); in CeedCompileCore_Cuda()
200 … CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "---------- COMPILE ERROR DETECTED ----------\n"); in CeedCompileCore_Cuda()
201 CeedDebug(ceed, "Error: %s\nCompile log:\n%s\n", nvrtcGetErrorString(result), log); in CeedCompileCore_Cuda()
202 CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "---------- BACKEND MAY FALLBACK ----------\n"); in CeedCompileCore_Cuda()
246 …(file, ceed, CEED_ERROR_BACKEND, "Failed to create file. Write access is required for cuda-clang"); in CeedCompileCore_Cuda()
260 …CeedDebug(ceed, "There are %d source dirs, including %s\n", num_rust_source_dirs, rust_source_dirs… in CeedCompileCore_Cuda()
276 // Compile Rust crate(s) needed in CeedCompileCore_Cuda()
280 …and = "cargo +" + std::string(rust_toolchain) + " build --release --target nvptx64-nvidia-cuda --c… in CeedCompileCore_Cuda()
281 "/.cargo/config.toml --manifest-path " + rust_dirs[i] + "/Cargo.toml"; in CeedCompileCore_Cuda()
286 bool use_llvm_version = ceed_data->use_llvm_version; in CeedCompileCore_Cuda()
287 int llvm_version = ceed_data->llvm_version; in CeedCompileCore_Cuda()
290 …nd $(rustup run " + std::string(rust_toolchain) + " rustc --print sysroot) -name llvm-link) --vers… in CeedCompileCore_Cuda()
291 CeedDebug(ceed, "Attempting to detect Rust LLVM version.\ncommand:\n$ %s", command.c_str()); in CeedCompileCore_Cuda()
302 CeedDebug(ceed, "output:\n%s", output.c_str()); in CeedCompileCore_Cuda()
303 …) == 0, ceed, CEED_ERROR_BACKEND, "Failed to detect Rust LLVM version\ncommand:\n$ %s\nerror:\n%s", in CeedCompileCore_Cuda()
314 ceed_data->llvm_version = llvm_version = std::stoi(version_substring); in CeedCompileCore_Cuda()
317 command = std::string("clang++-") + std::to_string(llvm_version); in CeedCompileCore_Cuda()
319 ceed_data->use_llvm_version = use_llvm_version = pclose(output_stream) == 0; in CeedCompileCore_Cuda()
321 ceed_data->llvm_version = -1; in CeedCompileCore_Cuda()
322 ceed_data->use_llvm_version = use_llvm_version = false; in CeedCompileCore_Cuda()
327 …ng++" + (use_llvm_version ? (std::string("-") + std::to_string(llvm_version)) : "") + " -flto=thin… in CeedCompileCore_Cuda()
328 …std::to_string(prop.major) + std::to_string(prop.minor) + " --cuda-device-only -emit-llvm -S temp/… in CeedCompileCore_Cuda()
329 "_0_source.cu -o temp/kernel_" + std::to_string(build_id) + "_1_wrapped.ll "; in CeedCompileCore_Cuda()
334 // Find Rust's llvm-link tool and run it in CeedCompileCore_Cuda()
335 …$(find $(rustup run " + std::string(rust_toolchain) + " rustc --print sysroot) -name llvm-link) te… in CeedCompileCore_Cuda()
337 "_1_wrapped.ll --ignore-non-bitcode --internalize --only-needed -S -o " in CeedCompileCore_Cuda()
345 std::string dir = rust_dirs[i] + "/target/nvptx64-nvidia-cuda/release"; in CeedCompileCore_Cuda()
348 … CeedCheck(dp != nullptr, ceed, CEED_ERROR_BACKEND, "Could not open directory: %s", dir.c_str()); in CeedCompileCore_Cuda()
353 std::string filename(entry->d_name); in CeedCompileCore_Cuda()
355 if (filename.size() >= 2 && filename.substr(filename.size() - 2) == ".a") { in CeedCompileCore_Cuda()
366 …" rustc --print sysroot) -name opt) --passes internalize,inline temp/kernel_" + std::to_string(bui… in CeedCompileCore_Cuda()
367 "_2_linked.ll -o temp/kernel_" + std::to_string(build_id) + "_3_opt.bc") in CeedCompileCore_Cuda()
372 …find $(rustup run " + std::string(rust_toolchain) + " rustc --print sysroot) -name llc) -O3 -mcpu=… in CeedCompileCore_Cuda()
374 "_3_opt.bc -o temp/kernel_" + std::to_string(build_id) + "_4_final.ptx") in CeedCompileCore_Cuda()
396 … CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "---------- COMPILE ERROR DETECTED ----------\n"); in CeedCompileCore_Cuda()
398 CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "---------- BACKEND MAY FALLBACK ----------\n"); in CeedCompileCore_Cuda()
430 //------------------------------------------------------------------------------
432 //------------------------------------------------------------------------------
438 //------------------------------------------------------------------------------
442 //------------------------------------------------------------------------------
451 //------------------------------------------------------------------------------
453 //------------------------------------------------------------------------------
459 //------------------------------------------------------------------------------
461 //------------------------------------------------------------------------------
468 //------------------------------------------------------------------------------
470 //------------------------------------------------------------------------------
491 CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "---------- LAUNCH ERROR DETECTED ----------\n"); in CeedRunKernelDimSharedCore_Cuda()
494 CeedDebug256(ceed, CEED_DEBUG_COLOR_WARNING, "---------- BACKEND MAY FALLBACK ----------\n"); in CeedRunKernelDimSharedCore_Cuda()
518 //------------------------------------------------------------------------------