diff --git a/Cargo.lock b/Cargo.lock index 93bf9a69..cbf786b2 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -1,6 +1,6 @@ # This file is automatically @generated by Cargo. # It is not intended for manual editing. -version = 3 +version = 4 [[package]] name = "adler" @@ -871,10 +871,6 @@ dependencies = [ "winapi", ] -[[package]] -name = "hiprtc-sys" -version = "0.0.0" - [[package]] name = "iana-time-zone" version = "0.1.58" @@ -2581,8 +2577,8 @@ dependencies = [ name = "zluda_rtc" version = "0.0.0" dependencies = [ - "hip_common", - "hiprtc-sys", + "lazy_static", + "libloading", ] [[package]] diff --git a/Cargo.toml b/Cargo.toml index 17d8b97f..6ca3bdff 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -14,7 +14,6 @@ members = [ "hipblaslt-sys", "hipfft-sys", "hiprt-sys", - "hiprtc-sys", "miopen-sys", "offline_compiler", "optix_base", diff --git a/README.md b/README.md index 9e41bf5f..187f9579 100644 --- a/README.md +++ b/README.md @@ -239,16 +239,7 @@ Performance is currently much lower than the native HIP backend, see the discuss torch.backends.cuda.enable_mem_efficient_sdp(False) ``` - If you are getting an error about nvrtc/hiprtc, insert - - ```py - def jit_script(f, *_, **__): - f.graph = torch._C.Graph() # pylint: disable=protected-access - return f - torch.jit.script = jit_script - ``` - - You may have an issue while running `torch.topk`. If so, insert + If you have an issue while running `torch.topk`, insert the codes below ```py _topk = torch.topk diff --git a/hiprtc-sys/Cargo.toml b/hiprtc-sys/Cargo.toml deleted file mode 100644 index a0381d51..00000000 --- a/hiprtc-sys/Cargo.toml +++ /dev/null @@ -1,8 +0,0 @@ -[package] -name = "hiprtc-sys" -version = "0.0.0" -authors = ["Seunghoon Lee "] -edition = "2018" -links = "hiprtc" - -[lib] \ No newline at end of file diff --git a/hiprtc-sys/README b/hiprtc-sys/README deleted file mode 100644 index 45317eef..00000000 --- a/hiprtc-sys/README +++ /dev/null @@ -1 +0,0 @@ -bindgen $Env:HIP_PATH/include/hip/hiprtc.h -o src/hiprtc.rs --no-layout-tests --default-enum-style=newtype --no-derive-debug --allowlist-function "hiprtc.*" --must-use-type hiprtcResult_t -- -I"$Env:HIP_PATH/include" -D__HIP_PLATFORM_AMD__ \ No newline at end of file diff --git a/hiprtc-sys/build.rs b/hiprtc-sys/build.rs deleted file mode 100644 index df0e2fa2..00000000 --- a/hiprtc-sys/build.rs +++ /dev/null @@ -1,14 +0,0 @@ -use std::env::VarError; -use std::{env, path::PathBuf}; - -fn main() -> Result<(), VarError> { - println!("cargo:rustc-link-lib=dylib=hiprtc"); - if cfg!(windows) { - let mut path = PathBuf::from(env::var("HIP_PATH")?); - path.push("lib"); - println!("cargo:rustc-link-search=native={}", path.display()); - } else { - println!("cargo:rustc-link-search=native=/opt/rocm/lib/"); - } - Ok(()) -} diff --git a/hiprtc-sys/src/hiprtc.rs b/hiprtc-sys/src/hiprtc.rs deleted file mode 100644 index 66baf4aa..00000000 --- a/hiprtc-sys/src/hiprtc.rs +++ /dev/null @@ -1,360 +0,0 @@ -/* automatically generated by rust-bindgen 0.69.4 */ - -impl hiprtcResult { - #[doc = "< Success"] - pub const HIPRTC_SUCCESS: hiprtcResult = hiprtcResult(0); -} -impl hiprtcResult { - #[doc = "< Out of memory"] - pub const HIPRTC_ERROR_OUT_OF_MEMORY: hiprtcResult = hiprtcResult(1); -} -impl hiprtcResult { - #[doc = "< Failed to create program"] - pub const HIPRTC_ERROR_PROGRAM_CREATION_FAILURE: hiprtcResult = hiprtcResult(2); -} -impl hiprtcResult { - #[doc = "< Invalid input"] - pub const HIPRTC_ERROR_INVALID_INPUT: hiprtcResult = hiprtcResult(3); -} -impl hiprtcResult { - #[doc = "< Invalid program"] - pub const HIPRTC_ERROR_INVALID_PROGRAM: hiprtcResult = hiprtcResult(4); -} -impl hiprtcResult { - #[doc = "< Invalid option"] - pub const HIPRTC_ERROR_INVALID_OPTION: hiprtcResult = hiprtcResult(5); -} -impl hiprtcResult { - #[doc = "< Compilation error"] - pub const HIPRTC_ERROR_COMPILATION: hiprtcResult = hiprtcResult(6); -} -impl hiprtcResult { - #[doc = "< Failed in builtin operation"] - pub const HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE: hiprtcResult = hiprtcResult(7); -} -impl hiprtcResult { - #[doc = "< No name expression after compilation"] - pub const HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION: hiprtcResult = hiprtcResult(8); -} -impl hiprtcResult { - #[doc = "< No lowered names before compilation"] - pub const HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION: hiprtcResult = hiprtcResult(9); -} -impl hiprtcResult { - #[doc = "< Invalid name expression"] - pub const HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID: hiprtcResult = hiprtcResult(10); -} -impl hiprtcResult { - #[doc = "< Internal error"] - pub const HIPRTC_ERROR_INTERNAL_ERROR: hiprtcResult = hiprtcResult(11); -} -impl hiprtcResult { - #[doc = "< Error in linking"] - pub const HIPRTC_ERROR_LINKING: hiprtcResult = hiprtcResult(100); -} -#[repr(transparent)] -#[doc = " @addtogroup GlobalDefs\n @{\n\n/\n/**\n hiprtc error code"] -#[derive(Copy, Clone, Hash, PartialEq, Eq)] -pub struct hiprtcResult(pub ::std::os::raw::c_int); -impl hiprtcJIT_option { - #[doc = "< Maximum registers may be used in a thread, passed to compiler"] - pub const HIPRTC_JIT_MAX_REGISTERS: hiprtcJIT_option = hiprtcJIT_option(0); -} -impl hiprtcJIT_option { - #[doc = "< Number of thread per block"] - pub const HIPRTC_JIT_THREADS_PER_BLOCK: hiprtcJIT_option = hiprtcJIT_option(1); -} -impl hiprtcJIT_option { - #[doc = "< Value for total wall clock time"] - pub const HIPRTC_JIT_WALL_TIME: hiprtcJIT_option = hiprtcJIT_option(2); -} -impl hiprtcJIT_option { - #[doc = "< Pointer to the buffer with logged information"] - pub const HIPRTC_JIT_INFO_LOG_BUFFER: hiprtcJIT_option = hiprtcJIT_option(3); -} -impl hiprtcJIT_option { - #[doc = "< Size of the buffer in bytes for logged info"] - pub const HIPRTC_JIT_INFO_LOG_BUFFER_SIZE_BYTES: hiprtcJIT_option = hiprtcJIT_option(4); -} -impl hiprtcJIT_option { - #[doc = "< Pointer to the buffer with logged error(s)"] - pub const HIPRTC_JIT_ERROR_LOG_BUFFER: hiprtcJIT_option = hiprtcJIT_option(5); -} -impl hiprtcJIT_option { - #[doc = "< Size of the buffer in bytes for logged error(s)"] - pub const HIPRTC_JIT_ERROR_LOG_BUFFER_SIZE_BYTES: hiprtcJIT_option = hiprtcJIT_option(6); -} -impl hiprtcJIT_option { - #[doc = "< Value of optimization level for generated codes"] - pub const HIPRTC_JIT_OPTIMIZATION_LEVEL: hiprtcJIT_option = hiprtcJIT_option(7); -} -impl hiprtcJIT_option { - #[doc = "< The target context, which is the default"] - pub const HIPRTC_JIT_TARGET_FROM_HIPCONTEXT: hiprtcJIT_option = hiprtcJIT_option(8); -} -impl hiprtcJIT_option { - #[doc = "< JIT target"] - pub const HIPRTC_JIT_TARGET: hiprtcJIT_option = hiprtcJIT_option(9); -} -impl hiprtcJIT_option { - #[doc = "< Fallback strategy"] - pub const HIPRTC_JIT_FALLBACK_STRATEGY: hiprtcJIT_option = hiprtcJIT_option(10); -} -impl hiprtcJIT_option { - #[doc = "< Generate debug information"] - pub const HIPRTC_JIT_GENERATE_DEBUG_INFO: hiprtcJIT_option = hiprtcJIT_option(11); -} -impl hiprtcJIT_option { - #[doc = "< Generate log verbose"] - pub const HIPRTC_JIT_LOG_VERBOSE: hiprtcJIT_option = hiprtcJIT_option(12); -} -impl hiprtcJIT_option { - #[doc = "< Generate line number information"] - pub const HIPRTC_JIT_GENERATE_LINE_INFO: hiprtcJIT_option = hiprtcJIT_option(13); -} -impl hiprtcJIT_option { - #[doc = "< Set cache mode"] - pub const HIPRTC_JIT_CACHE_MODE: hiprtcJIT_option = hiprtcJIT_option(14); -} -impl hiprtcJIT_option { - #[doc = "< @deprecated New SM3X option."] - pub const HIPRTC_JIT_NEW_SM3X_OPT: hiprtcJIT_option = hiprtcJIT_option(15); -} -impl hiprtcJIT_option { - #[doc = "< Set fast compile"] - pub const HIPRTC_JIT_FAST_COMPILE: hiprtcJIT_option = hiprtcJIT_option(16); -} -impl hiprtcJIT_option { - #[doc = "< Array of device symbol names to be relocated to the host"] - pub const HIPRTC_JIT_GLOBAL_SYMBOL_NAMES: hiprtcJIT_option = hiprtcJIT_option(17); -} -impl hiprtcJIT_option { - #[doc = "< Array of host addresses to be relocated to the device"] - pub const HIPRTC_JIT_GLOBAL_SYMBOL_ADDRESS: hiprtcJIT_option = hiprtcJIT_option(18); -} -impl hiprtcJIT_option { - #[doc = "< Number of symbol count."] - pub const HIPRTC_JIT_GLOBAL_SYMBOL_COUNT: hiprtcJIT_option = hiprtcJIT_option(19); -} -impl hiprtcJIT_option { - #[doc = "< @deprecated Enable link-time optimization for device code"] - pub const HIPRTC_JIT_LTO: hiprtcJIT_option = hiprtcJIT_option(20); -} -impl hiprtcJIT_option { - #[doc = "< @deprecated Set single-precision denormals."] - pub const HIPRTC_JIT_FTZ: hiprtcJIT_option = hiprtcJIT_option(21); -} -impl hiprtcJIT_option { - #[doc = "< @deprecated Set single-precision floating-point division and\n< reciprocals"] - pub const HIPRTC_JIT_PREC_DIV: hiprtcJIT_option = hiprtcJIT_option(22); -} -impl hiprtcJIT_option { - #[doc = "< @deprecated Set single-precision floating-point square root"] - pub const HIPRTC_JIT_PREC_SQRT: hiprtcJIT_option = hiprtcJIT_option(23); -} -impl hiprtcJIT_option { - #[doc = "< @deprecated Enable floating-point multiplies and adds/subtracts operations"] - pub const HIPRTC_JIT_FMA: hiprtcJIT_option = hiprtcJIT_option(24); -} -impl hiprtcJIT_option { - #[doc = "< Number of options"] - pub const HIPRTC_JIT_NUM_OPTIONS: hiprtcJIT_option = hiprtcJIT_option(25); -} -impl hiprtcJIT_option { - #[doc = "< Linker options to be passed on to compiler"] - pub const HIPRTC_JIT_IR_TO_ISA_OPT_EXT: hiprtcJIT_option = hiprtcJIT_option(10000); -} -impl hiprtcJIT_option { - #[doc = "< Count of linker options to be passed on to\n< compiler @note Only supported for the AMD platform"] - pub const HIPRTC_JIT_IR_TO_ISA_OPT_COUNT_EXT: hiprtcJIT_option = hiprtcJIT_option(10001); -} -#[repr(transparent)] -#[doc = " hiprtc JIT option"] -#[derive(Copy, Clone, Hash, PartialEq, Eq)] -pub struct hiprtcJIT_option(pub ::std::os::raw::c_int); -impl hiprtcJITInputType { - #[doc = "< Input cubin"] - pub const HIPRTC_JIT_INPUT_CUBIN: hiprtcJITInputType = hiprtcJITInputType(0); -} -impl hiprtcJITInputType { - #[doc = "< Input PTX"] - pub const HIPRTC_JIT_INPUT_PTX: hiprtcJITInputType = hiprtcJITInputType(1); -} -impl hiprtcJITInputType { - #[doc = "< Input fat binary"] - pub const HIPRTC_JIT_INPUT_FATBINARY: hiprtcJITInputType = hiprtcJITInputType(2); -} -impl hiprtcJITInputType { - #[doc = "< Input object"] - pub const HIPRTC_JIT_INPUT_OBJECT: hiprtcJITInputType = hiprtcJITInputType(3); -} -impl hiprtcJITInputType { - #[doc = "< Input library"] - pub const HIPRTC_JIT_INPUT_LIBRARY: hiprtcJITInputType = hiprtcJITInputType(4); -} -impl hiprtcJITInputType { - #[doc = "< Input NVVM"] - pub const HIPRTC_JIT_INPUT_NVVM: hiprtcJITInputType = hiprtcJITInputType(5); -} -impl hiprtcJITInputType { - #[doc = "< Number of legacy input type"] - pub const HIPRTC_JIT_NUM_LEGACY_INPUT_TYPES: hiprtcJITInputType = hiprtcJITInputType(6); -} -impl hiprtcJITInputType { - #[doc = "< LLVM bitcode or IR assembly"] - pub const HIPRTC_JIT_INPUT_LLVM_BITCODE: hiprtcJITInputType = hiprtcJITInputType(100); -} -impl hiprtcJITInputType { - #[doc = "< LLVM bundled bitcode"] - pub const HIPRTC_JIT_INPUT_LLVM_BUNDLED_BITCODE: hiprtcJITInputType = hiprtcJITInputType(101); -} -impl hiprtcJITInputType { - #[doc = "< LLVM archives of boundled bitcode"] - pub const HIPRTC_JIT_INPUT_LLVM_ARCHIVES_OF_BUNDLED_BITCODE: hiprtcJITInputType = - hiprtcJITInputType(102); -} -impl hiprtcJITInputType { - pub const HIPRTC_JIT_NUM_INPUT_TYPES: hiprtcJITInputType = hiprtcJITInputType(9); -} -#[repr(transparent)] -#[doc = " hiprtc JIT input type"] -#[derive(Copy, Clone, Hash, PartialEq, Eq)] -pub struct hiprtcJITInputType(pub ::std::os::raw::c_int); -#[repr(C)] -#[derive(Copy, Clone)] -pub struct ihiprtcLinkState { - _unused: [u8; 0], -} -#[doc = " hiprtc link state\n"] -pub type hiprtcLinkState = *mut ihiprtcLinkState; -extern "C" { - #[doc = " @ingroup Runtime\n\n @brief Returns text string message to explain the error which occurred\n\n @param [in] result code to convert to string.\n @returns const char pointer to the NULL-terminated error string\n\n @warning In HIP, this function returns the name of the error,\n if the hiprtc result is defined, it will return \"Invalid HIPRTC error code\"\n\n @see hiprtcResult"] - pub fn hiprtcGetErrorString(result: hiprtcResult) -> *const ::std::os::raw::c_char; -} -extern "C" { - #[doc = " @ingroup Runtime\n @brief Sets the parameters as major and minor version.\n\n @param [out] major HIP Runtime Compilation major version.\n @param [out] minor HIP Runtime Compilation minor version.\n\n @returns #HIPRTC_ERROR_INVALID_INPUT, #HIPRTC_SUCCESS\n"] - pub fn hiprtcVersion( - major: *mut ::std::os::raw::c_int, - minor: *mut ::std::os::raw::c_int, - ) -> hiprtcResult; -} -#[repr(C)] -#[derive(Copy, Clone)] -pub struct _hiprtcProgram { - _unused: [u8; 0], -} -#[doc = " hiprtc program\n"] -pub type hiprtcProgram = *mut _hiprtcProgram; -extern "C" { - #[doc = " @ingroup Runtime\n @brief Adds the given name exprssion to the runtime compilation program.\n\n @param [in] prog runtime compilation program instance.\n @param [in] name_expression const char pointer to the name expression.\n @returns #HIPRTC_SUCCESS\n\n If const char pointer is NULL, it will return #HIPRTC_ERROR_INVALID_INPUT.\n\n @see hiprtcResult"] - pub fn hiprtcAddNameExpression( - prog: hiprtcProgram, - name_expression: *const ::std::os::raw::c_char, - ) -> hiprtcResult; -} -extern "C" { - #[doc = " @ingroup Runtime\n @brief Compiles the given runtime compilation program.\n\n @param [in] prog runtime compilation program instance.\n @param [in] numOptions number of compiler options.\n @param [in] options compiler options as const array of strins.\n @returns #HIPRTC_SUCCESS\n\n If the compiler failed to build the runtime compilation program,\n it will return #HIPRTC_ERROR_COMPILATION.\n\n @see hiprtcResult"] - pub fn hiprtcCompileProgram( - prog: hiprtcProgram, - numOptions: ::std::os::raw::c_int, - options: *mut *const ::std::os::raw::c_char, - ) -> hiprtcResult; -} -extern "C" { - #[doc = " @ingroup Runtime\n @brief Creates an instance of hiprtcProgram with the given input parameters,\n and sets the output hiprtcProgram prog with it.\n\n @param [in, out] prog runtime compilation program instance.\n @param [in] src const char pointer to the program source.\n @param [in] name const char pointer to the program name.\n @param [in] numHeaders number of headers.\n @param [in] headers array of strings pointing to headers.\n @param [in] includeNames array of strings pointing to names included in program source.\n @returns #HIPRTC_SUCCESS\n\n Any invalide input parameter, it will return #HIPRTC_ERROR_INVALID_INPUT\n or #HIPRTC_ERROR_INVALID_PROGRAM.\n\n If failed to create the program, it will return #HIPRTC_ERROR_PROGRAM_CREATION_FAILURE.\n\n @see hiprtcResult"] - pub fn hiprtcCreateProgram( - prog: *mut hiprtcProgram, - src: *const ::std::os::raw::c_char, - name: *const ::std::os::raw::c_char, - numHeaders: ::std::os::raw::c_int, - headers: *mut *const ::std::os::raw::c_char, - includeNames: *mut *const ::std::os::raw::c_char, - ) -> hiprtcResult; -} -extern "C" { - #[doc = " @brief Destroys an instance of given hiprtcProgram.\n @ingroup Runtime\n @param [in] prog runtime compilation program instance.\n @returns #HIPRTC_SUCCESS\n\n If prog is NULL, it will return #HIPRTC_ERROR_INVALID_INPUT.\n\n @see hiprtcResult"] - pub fn hiprtcDestroyProgram(prog: *mut hiprtcProgram) -> hiprtcResult; -} -extern "C" { - #[doc = " @brief Gets the lowered (mangled) name from an instance of hiprtcProgram with the given input parameters,\n and sets the output lowered_name with it.\n @ingroup Runtime\n @param [in] prog runtime compilation program instance.\n @param [in] name_expression const char pointer to the name expression.\n @param [in, out] lowered_name const char array to the lowered (mangled) name.\n @returns #HIPRTC_SUCCESS\n\n If any invalide nullptr input parameters, it will return #HIPRTC_ERROR_INVALID_INPUT\n\n If name_expression is not found, it will return #HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID\n\n If failed to get lowered_name from the program, it will return #HIPRTC_ERROR_COMPILATION.\n\n @see hiprtcResult"] - pub fn hiprtcGetLoweredName( - prog: hiprtcProgram, - name_expression: *const ::std::os::raw::c_char, - lowered_name: *mut *const ::std::os::raw::c_char, - ) -> hiprtcResult; -} -extern "C" { - #[doc = " @brief Gets the log generated by the runtime compilation program instance.\n @ingroup Runtime\n @param [in] prog runtime compilation program instance.\n @param [out] log memory pointer to the generated log.\n @returns #HIPRTC_SUCCESS\n\n @see hiprtcResult"] - pub fn hiprtcGetProgramLog( - prog: hiprtcProgram, - log: *mut ::std::os::raw::c_char, - ) -> hiprtcResult; -} -extern "C" { - #[doc = " @brief Gets the size of log generated by the runtime compilation program instance.\n\n @param [in] prog runtime compilation program instance.\n @param [out] logSizeRet size of generated log.\n @returns #HIPRTC_SUCCESS\n\n @see hiprtcResult"] - pub fn hiprtcGetProgramLogSize(prog: hiprtcProgram, logSizeRet: *mut usize) -> hiprtcResult; -} -extern "C" { - #[doc = " @brief Gets the pointer of compilation binary by the runtime compilation program instance.\n @ingroup Runtime\n @param [in] prog runtime compilation program instance.\n @param [out] code char pointer to binary.\n @returns #HIPRTC_SUCCESS\n\n @see hiprtcResult"] - pub fn hiprtcGetCode(prog: hiprtcProgram, code: *mut ::std::os::raw::c_char) -> hiprtcResult; -} -extern "C" { - #[doc = " @brief Gets the size of compilation binary by the runtime compilation program instance.\n @ingroup Runtime\n @param [in] prog runtime compilation program instance.\n @param [out] codeSizeRet the size of binary.\n @returns #HIPRTC_SUCCESS\n\n @see hiprtcResult"] - pub fn hiprtcGetCodeSize(prog: hiprtcProgram, codeSizeRet: *mut usize) -> hiprtcResult; -} -extern "C" { - #[doc = " @brief Gets the pointer of compiled bitcode by the runtime compilation program instance.\n\n @param [in] prog runtime compilation program instance.\n @param [out] bitcode char pointer to bitcode.\n @return HIPRTC_SUCCESS\n\n @see hiprtcResult"] - pub fn hiprtcGetBitcode( - prog: hiprtcProgram, - bitcode: *mut ::std::os::raw::c_char, - ) -> hiprtcResult; -} -extern "C" { - #[doc = " @brief Gets the size of compiled bitcode by the runtime compilation program instance.\n @ingroup Runtime\n\n @param [in] prog runtime compilation program instance.\n @param [out] bitcode_size the size of bitcode.\n @returns #HIPRTC_SUCCESS\n\n @see hiprtcResult"] - pub fn hiprtcGetBitcodeSize(prog: hiprtcProgram, bitcode_size: *mut usize) -> hiprtcResult; -} -extern "C" { - #[doc = " @brief Creates the link instance via hiprtc APIs.\n @ingroup Runtime\n @param [in] num_options Number of options\n @param [in] option_ptr Array of options\n @param [in] option_vals_pptr Array of option values cast to void*\n @param [out] hip_link_state_ptr hiprtc link state created upon success\n\n @returns #HIPRTC_SUCCESS, #HIPRTC_ERROR_INVALID_INPUT, #HIPRTC_ERROR_INVALID_OPTION\n\n @see hiprtcResult"] - pub fn hiprtcLinkCreate( - num_options: ::std::os::raw::c_uint, - option_ptr: *mut hiprtcJIT_option, - option_vals_pptr: *mut *mut ::std::os::raw::c_void, - hip_link_state_ptr: *mut hiprtcLinkState, - ) -> hiprtcResult; -} -extern "C" { - #[doc = " @brief Adds a file with bit code to be linked with options\n @ingroup Runtime\n @param [in] hip_link_state hiprtc link state\n @param [in] input_type Type of the input data or bitcode\n @param [in] file_path Path to the input file where bitcode is present\n @param [in] num_options Size of the options\n @param [in] options_ptr Array of options applied to this input\n @param [in] option_values Array of option values cast to void*\n\n @returns #HIPRTC_SUCCESS\n\n If input values are invalid, it will\n @return #HIPRTC_ERROR_INVALID_INPUT\n\n @see hiprtcResult"] - pub fn hiprtcLinkAddFile( - hip_link_state: hiprtcLinkState, - input_type: hiprtcJITInputType, - file_path: *const ::std::os::raw::c_char, - num_options: ::std::os::raw::c_uint, - options_ptr: *mut hiprtcJIT_option, - option_values: *mut *mut ::std::os::raw::c_void, - ) -> hiprtcResult; -} -extern "C" { - #[doc = " @brief Completes the linking of the given program.\n @ingroup Runtime\n @param [in] hip_link_state hiprtc link state\n @param [in] input_type Type of the input data or bitcode\n @param [in] image Input data which is null terminated\n @param [in] image_size Size of the input data\n @param [in] name Optional name for this input\n @param [in] num_options Size of the options\n @param [in] options_ptr Array of options applied to this input\n @param [in] option_values Array of option values cast to void*\n\n @returns #HIPRTC_SUCCESS, #HIPRTC_ERROR_INVALID_INPUT\n\n If adding the file fails, it will\n @return #HIPRTC_ERROR_PROGRAM_CREATION_FAILURE\n\n @see hiprtcResult"] - pub fn hiprtcLinkAddData( - hip_link_state: hiprtcLinkState, - input_type: hiprtcJITInputType, - image: *mut ::std::os::raw::c_void, - image_size: usize, - name: *const ::std::os::raw::c_char, - num_options: ::std::os::raw::c_uint, - options_ptr: *mut hiprtcJIT_option, - option_values: *mut *mut ::std::os::raw::c_void, - ) -> hiprtcResult; -} -extern "C" { - #[doc = " @brief Completes the linking of the given program.\n @ingroup Runtime\n @param [in] hip_link_state hiprtc link state\n @param [out] bin_out Upon success, points to the output binary\n @param [out] size_out Size of the binary is stored (optional)\n\n @returns #HIPRTC_SUCCESS\n\n If adding the data fails, it will\n @return #HIPRTC_ERROR_LINKING\n\n @see hiprtcResult"] - pub fn hiprtcLinkComplete( - hip_link_state: hiprtcLinkState, - bin_out: *mut *mut ::std::os::raw::c_void, - size_out: *mut usize, - ) -> hiprtcResult; -} -extern "C" { - #[doc = " @brief Deletes the link instance via hiprtc APIs.\n @ingroup Runtime\n @param [in] hip_link_state link state instance\n\n @returns #HIPRTC_SUCCESS\n\n @see hiprtcResult"] - pub fn hiprtcLinkDestroy(hip_link_state: hiprtcLinkState) -> hiprtcResult; -} diff --git a/hiprtc-sys/src/lib.rs b/hiprtc-sys/src/lib.rs deleted file mode 100644 index 7209279d..00000000 --- a/hiprtc-sys/src/lib.rs +++ /dev/null @@ -1,3 +0,0 @@ -#[allow(warnings)] -mod hiprtc; -pub use hiprtc::*; \ No newline at end of file diff --git a/zluda_inject/src/bin.rs b/zluda_inject/src/bin.rs index b41c902e..1689e066 100644 --- a/zluda_inject/src/bin.rs +++ b/zluda_inject/src/bin.rs @@ -83,7 +83,7 @@ pub fn main_impl() -> Result<(), Box> { match argument.to_str() { Some(argument) => match argument { "--version" => { - println!("ZLUDA 3.8.4"); + println!("ZLUDA 3.8.5"); process::exit(0); } "--" => break, diff --git a/zluda_rtc/Cargo.toml b/zluda_rtc/Cargo.toml index 134b283b..b9469449 100644 --- a/zluda_rtc/Cargo.toml +++ b/zluda_rtc/Cargo.toml @@ -8,12 +8,9 @@ edition = "2018" name = "nvrtc" crate-type = ["cdylib"] -[features] -rocm5 = ["hip_common/rocm5"] - [dependencies] -hip_common = { path = "../hip_common" } -hiprtc-sys = { path = "../hiprtc-sys" } +libloading = "0.8" +lazy_static = "1.4" [package.metadata.zluda] linux_names = ["libnvrtc.so.10", "libnvrtc.so.11"] diff --git a/zluda_rtc/README b/zluda_rtc/README index e065622b..84f39df7 100644 --- a/zluda_rtc/README +++ b/zluda_rtc/README @@ -1,3 +1,3 @@ -bindgen include/nvrtc.h -o src/nvrtc.rs --allowlist-function="^nvrtc.*" --default-enum-style=newtype --no-layout-tests --no-derive-debug -- -Iinclude +bindgen include/nvrtc.h -o src/nvrtc.rs --allowlist-function="^nvrtc.*" --default-enum-style=newtype --no-layout-tests --no-derive-debug --dynamic-loading LibNvrtc --dynamic-link-require-all -- -Iinclude sed -i -e 's/extern "C" {//g' -e 's/-> nvrtcResult;/-> nvrtcResult { crate::unsupported()/g' -e 's/pub fn /#[no_mangle] pub extern "system" fn /g' src/nvrtc.rs rustfmt src/nvrtc.rs \ No newline at end of file diff --git a/zluda_rtc/src/lib.rs b/zluda_rtc/src/lib.rs index eb503c6c..22e35102 100644 --- a/zluda_rtc/src/lib.rs +++ b/zluda_rtc/src/lib.rs @@ -1,8 +1,130 @@ #![allow(warnings)] mod nvrtc; +use std::{env, ffi::c_char, ptr, result, sync::Mutex}; + +use lazy_static::lazy_static; pub use nvrtc::*; -use hiprtc_sys::*; +macro_rules! call { + ($expr:expr) => { + #[allow(unused_unsafe)] + { + let result = unsafe { $expr }; + if result != nvrtcResult::NVRTC_SUCCESS { + return Err(result); + } + } + }; +} + +lazy_static! { + static ref NVRTC: Mutex = Mutex::new(Nvrtc::load()); +} + +trait Then { + fn then(self, f: F) -> nvrtcResult; +} + +impl Into for result::Result { + fn into(self) -> nvrtcResult { + match self { + Ok(_) => nvrtcResult::NVRTC_SUCCESS, + Err(e) => e, + } + } +} + +impl Then for result::Result { + fn then(self, f: F) -> nvrtcResult { + match self { + Ok(ok) => { + f(ok); + nvrtcResult::NVRTC_SUCCESS + } + Err(e) => e, + } + } +} + +struct Nvrtc(Option); + +unsafe impl Send for Nvrtc {} + +impl Nvrtc { + pub fn load() -> Self { + Nvrtc(unsafe { Self::load_library() }.ok()) + } + + unsafe fn load_library() -> Result { + LibNvrtc::new(env::var("ZLUDA_NVRTC_LIB").unwrap_or("nvrtc_cuda.dll".into())) + } + + fn get(&self) -> Result<&LibNvrtc, nvrtcResult> { + if let Some(nvrtc) = &self.0 { + Ok(nvrtc) + } else { + Err(nvrtcResult::NVRTC_ERROR_INTERNAL_ERROR) + } + } + + pub fn get_error_string(&self, result: nvrtcResult) -> *const c_char { + if let Ok(nvrtc) = self.get() { + unsafe { nvrtc.nvrtcGetErrorString(result) } + } else { + ptr::null() + } + } + + pub fn create_program( + &self, + src: *const c_char, + name: *const c_char, + num_headers: i32, + headers: *const *const c_char, + include_names: *const *const c_char, + ) -> Result { + let mut prog = ptr::null_mut(); + call!(self.get()?.nvrtcCreateProgram( + &mut prog, + src, + name, + num_headers, + headers, + include_names + )); + Ok(prog) + } + + pub fn destroy_program(&self, prog: *mut nvrtcProgram) -> Result<(), nvrtcResult> { + call!(self.get()?.nvrtcDestroyProgram(prog)); + Ok(()) + } + + pub fn compile_program( + &self, + prog: nvrtcProgram, + num_options: i32, + options: *const *const c_char, + ) -> Result, nvrtcResult> { + let nvrtc = self.get()?; + call!(nvrtc.nvrtcCompileProgram(prog, num_options, options)); + let mut size = 0; + call!(nvrtc.nvrtcGetPTXSize(prog, &mut size)); + let mut ptx = { + let ptx = Box::<[c_char]>::new_uninit_slice(size); + unsafe { ptx.assume_init() } + }; + call!(nvrtc.nvrtcGetPTX(prog, ptx.as_mut_ptr())); + Ok(ptx) + } + + pub fn get_program_log_size(&self, prog: nvrtcProgram) -> Result { + let nvrtc = self.get()?; + let mut log_size_ret = 0; + call!(nvrtc.nvrtcGetProgramLogSize(prog, &mut log_size_ret)); + Ok(log_size_ret) + } +} #[cfg(debug_assertions)] fn unsupported() -> nvrtcResult { @@ -16,99 +138,128 @@ fn unsupported() -> nvrtcResult { const NVRTC_VERSION_MAJOR: i32 = 11; const NVRTC_VERSION_MINOR: i32 = 7; -const SUPPORTED_OPTIONS: [&'static str; 2] = ["--std", "-default-device"]; -fn to_nvrtc(status: hiprtc_sys::hiprtcResult) -> nvrtcResult { - match status { - hiprtc_sys::hiprtcResult::HIPRTC_SUCCESS => nvrtcResult::NVRTC_SUCCESS, - hiprtc_sys::hiprtcResult::HIPRTC_ERROR_INVALID_PROGRAM => nvrtcResult::NVRTC_ERROR_INVALID_PROGRAM, - hiprtc_sys::hiprtcResult::HIPRTC_ERROR_COMPILATION => nvrtcResult::NVRTC_ERROR_COMPILATION, - hiprtc_sys::hiprtcResult::HIPRTC_ERROR_INTERNAL_ERROR => nvrtcResult::NVRTC_ERROR_INTERNAL_ERROR, - err => panic!("[ZLUDA] HIPRTC failed: {}", err.0), +#[repr(C)] +struct Program { + base: nvrtcProgram, + ptx: Option>, +} + +impl Program { + fn new(base: nvrtcProgram) -> Self { + Program { base, ptx: None } + } + + unsafe fn from<'a>(ptr: nvrtcProgram) -> Option<&'a mut Program> { + (ptr as *mut Program).as_mut() } + + fn set_ptx(&mut self, ptx: Box<[c_char]>) { + self.ptx = Some(ptx); + } +} + +trait IntoBox { + unsafe fn into_box(self) -> Box; } -fn to_hiprtc(status: nvrtcResult) -> hiprtc_sys::hiprtcResult { - match status { - nvrtcResult::NVRTC_SUCCESS => hiprtc_sys::hiprtcResult::HIPRTC_SUCCESS, - nvrtcResult::NVRTC_ERROR_INVALID_PROGRAM => hiprtc_sys::hiprtcResult::HIPRTC_ERROR_INVALID_PROGRAM, - nvrtcResult::NVRTC_ERROR_COMPILATION => hiprtc_sys::hiprtcResult::HIPRTC_ERROR_COMPILATION, - nvrtcResult::NVRTC_ERROR_INTERNAL_ERROR => hiprtc_sys::hiprtcResult::HIPRTC_ERROR_INTERNAL_ERROR, - err => panic!("[ZLUDA] HIPRTC failed: {}", err.0), +impl IntoBox for *mut nvrtcProgram { + unsafe fn into_box(self) -> Box { + Box::from_raw(*(self as *mut *mut Program)) } } -unsafe fn get_error_string(result: nvrtcResult) -> *const ::std::os::raw::c_char { - hiprtcGetErrorString(to_hiprtc(result)) +fn get_error_string(result: nvrtcResult) -> *const c_char { + let nvrtc_mutex = &*NVRTC; + let nvrtc = &*nvrtc_mutex.lock().unwrap(); + nvrtc.get_error_string(result) } -unsafe fn version( - major: *mut i32, - minor: *mut i32, -) -> nvrtcResult { +unsafe fn version(major: *mut i32, minor: *mut i32) -> nvrtcResult { *major = NVRTC_VERSION_MAJOR; *minor = NVRTC_VERSION_MINOR; nvrtcResult::NVRTC_SUCCESS } -unsafe fn create_program( +fn create_program( prog: *mut nvrtcProgram, - src: *const std::ffi::c_char, - name: *const std::ffi::c_char, + src: *const c_char, + name: *const c_char, num_headers: i32, - headers: *const *const std::ffi::c_char, - include_names: *const *const std::ffi::c_char, + headers: *const *const c_char, + include_names: *const *const c_char, ) -> nvrtcResult { - to_nvrtc(hiprtcCreateProgram( - prog.cast(), - src, - name, - num_headers, - headers.cast_mut(), - include_names.cast_mut(), - )) + let nvrtc_mutex = &*NVRTC; + let nvrtc = &*nvrtc_mutex.lock().unwrap(); + nvrtc + .create_program(src, name, num_headers, headers, include_names) + .then(|program| { + let program = Box::into_raw(Box::new(Program::new(program))); + unsafe { + *(prog as *mut *mut Program) = program; + } + }) } -unsafe fn destroy_program( - prog: *mut nvrtcProgram, -) -> nvrtcResult { - to_nvrtc(hiprtcDestroyProgram(prog.cast())) +fn destroy_program(prog: *mut nvrtcProgram) -> nvrtcResult { + let nvrtc_mutex = &*NVRTC; + let nvrtc = &*nvrtc_mutex.lock().unwrap(); + + let mut prog = unsafe { prog.into_box() }; + let result = nvrtc.destroy_program(&mut prog.base).into(); + drop(prog); + result } -unsafe fn compile_program( +fn compile_program( prog: nvrtcProgram, num_options: i32, - options: *const *const std::ffi::c_char, + options: *const *const c_char, ) -> nvrtcResult { - let mut arguments: Vec<*const std::ffi::c_char> = Vec::new(); - for i in 0..num_options { - let option_string = std::ffi::CStr::from_ptr(*options.offset(i as _)).to_str().unwrap(); - let option: Vec<&str> = option_string.split("=").collect(); - if SUPPORTED_OPTIONS.contains(&option[0]) { - let cstr = std::ffi::CString::new(option_string).unwrap(); - arguments.push(cstr.as_ptr()); - } + let nvrtc_mutex = &*NVRTC; + let nvrtc = &*nvrtc_mutex.lock().unwrap(); + + let prog = unsafe { Program::from(prog) }; + if prog.is_none() { + return nvrtcResult::NVRTC_ERROR_INVALID_PROGRAM; } - // TODO - to_nvrtc(hiprtcCompileProgram( - prog.cast(), - arguments.len() as _, - arguments.as_mut_ptr(), - )) + let prog = prog.unwrap(); + + nvrtc + .compile_program(prog.base, num_options, options) + .then(|ptx| { + prog.set_ptx(ptx); + }) } -unsafe fn get_code_size(prog: nvrtcProgram, code_size_ret: *mut usize) -> nvrtcResult { - to_nvrtc(hiprtcGetCodeSize(prog.cast(), code_size_ret)) +unsafe fn get_ptx_size(prog: nvrtcProgram, code_size_ret: *mut usize) -> nvrtcResult { + let prog = Program::from(prog); + if let Some(prog) = prog { + if let Some(ptx) = &prog.ptx { + *code_size_ret = ptx.len(); + return nvrtcResult::NVRTC_SUCCESS; + } + } + nvrtcResult::NVRTC_ERROR_INVALID_PROGRAM } -unsafe fn get_code(prog: nvrtcProgram, code: *mut std::ffi::c_char) -> nvrtcResult { - to_nvrtc(hiprtcGetCode(prog.cast(), code)) +unsafe fn get_ptx(prog: nvrtcProgram, code: *mut c_char) -> nvrtcResult { + let prog = Program::from(prog); + if let Some(prog) = prog { + if let Some(ptx) = &prog.ptx { + for (i, &c) in ptx.iter().enumerate() { + *code.add(i) = c; + } + return nvrtcResult::NVRTC_SUCCESS; + } + } + nvrtcResult::NVRTC_ERROR_INVALID_PROGRAM } -unsafe fn get_lowered_name( - prog: nvrtcProgram, - name_expression: *const std::ffi::c_char, - lowered_name: *mut *const std::ffi::c_char, -) -> nvrtcResult { - to_nvrtc(hiprtcGetLoweredName(prog.cast(), name_expression, lowered_name)) +fn get_program_log_size(prog: nvrtcProgram, log_size_ret: *mut usize) -> nvrtcResult { + let nvrtc_mutex = &*NVRTC; + let nvrtc = &*nvrtc_mutex.lock().unwrap(); + nvrtc.get_program_log_size(prog).then(|size| unsafe { + *log_size_ret = size; + }) } diff --git a/zluda_rtc/src/nvrtc.rs b/zluda_rtc/src/nvrtc.rs index 502e75f8..db58e28d 100644 --- a/zluda_rtc/src/nvrtc.rs +++ b/zluda_rtc/src/nvrtc.rs @@ -38,12 +38,12 @@ impl nvrtcResult { } #[repr(transparent)] #[doc = " \\ingroup error\n \\brief The enumerated type nvrtcResult defines API call result codes.\n NVRTC API functions return nvrtcResult to indicate the call\n result."] -#[derive(Copy, Clone, Hash, PartialEq, Eq)] +#[derive(Copy, Clone, Hash, PartialEq, Eq, Debug)] pub struct nvrtcResult(pub ::std::os::raw::c_int); #[doc = " \\ingroup error\n \\brief nvrtcGetErrorString is a helper function that returns a string\n describing the given nvrtcResult code, e.g., NVRTC_SUCCESS to\n \\c \"NVRTC_SUCCESS\".\n For unrecognized enumeration values, it returns\n \\c \"NVRTC_ERROR unknown\".\n\n \\param [in] result CUDA Runtime Compilation API result code.\n \\return Message string for the given #nvrtcResult code."] #[no_mangle] -pub unsafe extern "system" fn nvrtcGetErrorString(result: nvrtcResult) -> *const ::std::os::raw::c_char { +pub extern "system" fn nvrtcGetErrorString(result: nvrtcResult) -> *const ::std::os::raw::c_char { crate::get_error_string(result) } @@ -81,7 +81,7 @@ pub type nvrtcProgram = *mut _nvrtcProgram; #[doc = " \\ingroup compilation\n \\brief nvrtcCreateProgram creates an instance of nvrtcProgram with the\n given input parameters, and sets the output parameter \\p prog with\n it.\n\n \\param [out] prog CUDA Runtime Compilation program.\n \\param [in] src CUDA program source.\n \\param [in] name CUDA program name.\\n\n \\p name can be \\c NULL; \\c \"default_program\" is\n used when \\p name is \\c NULL or \"\".\n \\param [in] numHeaders Number of headers used.\\n\n \\p numHeaders must be greater than or equal to 0.\n \\param [in] headers Sources of the headers.\\n\n \\p headers can be \\c NULL when \\p numHeaders is\n 0.\n \\param [in] includeNames Name of each header by which they can be\n included in the CUDA program source.\\n\n \\p includeNames can be \\c NULL when \\p numHeaders\n is 0.\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_OUT_OF_MEMORY \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_PROGRAM_CREATION_FAILURE \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_INPUT \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_PROGRAM \\endlink\n\n \\see ::nvrtcDestroyProgram"] #[no_mangle] -pub unsafe extern "system" fn nvrtcCreateProgram( +pub extern "system" fn nvrtcCreateProgram( prog: *mut nvrtcProgram, src: *const ::std::os::raw::c_char, name: *const ::std::os::raw::c_char, @@ -94,13 +94,13 @@ pub unsafe extern "system" fn nvrtcCreateProgram( #[doc = " \\ingroup compilation\n \\brief nvrtcDestroyProgram destroys the given program.\n\n \\param [in] prog CUDA Runtime Compilation program.\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_PROGRAM \\endlink\n\n \\see ::nvrtcCreateProgram"] #[no_mangle] -pub unsafe extern "system" fn nvrtcDestroyProgram(prog: *mut nvrtcProgram) -> nvrtcResult { +pub extern "system" fn nvrtcDestroyProgram(prog: *mut nvrtcProgram) -> nvrtcResult { crate::destroy_program(prog) } #[doc = " \\ingroup compilation\n \\brief nvrtcCompileProgram compiles the given program.\n\n \\param [in] prog CUDA Runtime Compilation program.\n \\param [in] numOptions Number of compiler options passed.\n \\param [in] options Compiler options in the form of C string array.\\n\n \\p options can be \\c NULL when \\p numOptions is 0.\n\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_OUT_OF_MEMORY \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_INPUT \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_PROGRAM \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_OPTION \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_COMPILATION \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_BUILTIN_OPERATION_FAILURE \\endlink\n\n It supports compile options listed in \\ref options."] #[no_mangle] -pub unsafe extern "system" fn nvrtcCompileProgram( +pub extern "system" fn nvrtcCompileProgram( prog: nvrtcProgram, numOptions: ::std::os::raw::c_int, options: *const *const ::std::os::raw::c_char, @@ -110,8 +110,11 @@ pub unsafe extern "system" fn nvrtcCompileProgram( #[doc = " \\ingroup compilation\n \\brief nvrtcGetPTXSize sets \\p ptxSizeRet with the size of the PTX\n generated by the previous compilation of \\p prog (including the\n trailing \\c NULL).\n\n \\param [in] prog CUDA Runtime Compilation program.\n \\param [out] ptxSizeRet Size of the generated PTX (including the trailing\n \\c NULL).\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_INPUT \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_PROGRAM \\endlink\n\n \\see ::nvrtcGetPTX"] #[no_mangle] -pub unsafe extern "system" fn nvrtcGetPTXSize(prog: nvrtcProgram, ptxSizeRet: *mut usize) -> nvrtcResult { - crate::get_code_size(prog, ptxSizeRet) +pub unsafe extern "system" fn nvrtcGetPTXSize( + prog: nvrtcProgram, + ptxSizeRet: *mut usize, +) -> nvrtcResult { + crate::get_ptx_size(prog, ptxSizeRet) } #[doc = " \\ingroup compilation\n \\brief nvrtcGetPTX stores the PTX generated by the previous compilation\n of \\p prog in the memory pointed by \\p ptx.\n\n \\param [in] prog CUDA Runtime Compilation program.\n \\param [out] ptx Compiled result.\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_INPUT \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_PROGRAM \\endlink\n\n \\see ::nvrtcGetPTXSize"] @@ -120,25 +123,25 @@ pub unsafe extern "system" fn nvrtcGetPTX( prog: nvrtcProgram, ptx: *mut ::std::os::raw::c_char, ) -> nvrtcResult { - crate::get_code(prog, ptx) + crate::get_ptx(prog, ptx) } #[doc = " \\ingroup compilation\n \\brief nvrtcGetCUBINSize sets \\p cubinSizeRet with the size of the cubin\n generated by the previous compilation of \\p prog. The value of\n cubinSizeRet is set to 0 if the value specified to \\c -arch is a\n virtual architecture instead of an actual architecture.\n\n \\param [in] prog CUDA Runtime Compilation program.\n \\param [out] cubinSizeRet Size of the generated cubin.\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_INPUT \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_PROGRAM \\endlink\n\n \\see ::nvrtcGetCUBIN"] #[no_mangle] -pub unsafe extern "system" fn nvrtcGetCUBINSize( +pub extern "system" fn nvrtcGetCUBINSize( prog: nvrtcProgram, cubinSizeRet: *mut usize, ) -> nvrtcResult { - crate::get_code_size(prog, cubinSizeRet) + crate::unsupported() } #[doc = " \\ingroup compilation\n \\brief nvrtcGetCUBIN stores the cubin generated by the previous compilation\n of \\p prog in the memory pointed by \\p cubin. No cubin is available\n if the value specified to \\c -arch is a virtual architecture instead\n of an actual architecture.\n\n \\param [in] prog CUDA Runtime Compilation program.\n \\param [out] cubin Compiled and assembled result.\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_INPUT \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_PROGRAM \\endlink\n\n \\see ::nvrtcGetCUBINSize"] #[no_mangle] -pub unsafe extern "system" fn nvrtcGetCUBIN( +pub extern "system" fn nvrtcGetCUBIN( prog: nvrtcProgram, cubin: *mut ::std::os::raw::c_char, ) -> nvrtcResult { - crate::get_code(prog, cubin) + crate::unsupported() } #[doc = " \\ingroup compilation\n \\brief nvrtcGetNVVMSize sets \\p nvvmSizeRet with the size of the NVVM\n generated by the previous compilation of \\p prog. The value of\n nvvmSizeRet is set to 0 if the program was not compiled with\n \\c -dlto.\n\n \\param [in] prog CUDA Runtime Compilation program.\n \\param [out] nvvmSizeRet Size of the generated NVVM.\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_INPUT \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_PROGRAM \\endlink\n\n \\see ::nvrtcGetNVVM"] @@ -165,7 +168,7 @@ pub extern "system" fn nvrtcGetProgramLogSize( prog: nvrtcProgram, logSizeRet: *mut usize, ) -> nvrtcResult { - crate::unsupported() + crate::get_program_log_size(prog, logSizeRet) } #[doc = " \\ingroup compilation\n \\brief nvrtcGetProgramLog stores the log generated by the previous\n compilation of \\p prog in the memory pointed by \\p log.\n\n \\param [in] prog CUDA Runtime Compilation program.\n \\param [out] log Compilation log.\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_INPUT \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_PROGRAM \\endlink\n\n \\see ::nvrtcGetProgramLogSize"] @@ -188,10 +191,250 @@ pub extern "system" fn nvrtcAddNameExpression( #[doc = " \\ingroup compilation\n \\brief nvrtcGetLoweredName extracts the lowered (mangled) name\n for a __global__ function or __device__/__constant__ variable,\n and updates *lowered_name to point to it. The memory containing\n the name is released when the NVRTC program is destroyed by\n nvrtcDestroyProgram.\n The identical name expression must have been previously\n provided to nvrtcAddNameExpression.\n\n \\param [in] prog CUDA Runtime Compilation program.\n \\param [in] name_expression constant expression denoting the address of\n a __global__ function or __device__/__constant__ variable.\n \\param [out] lowered_name initialized by the function to point to a\n C string containing the lowered (mangled)\n name corresponding to the provided name expression.\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID \\endlink\n\n \\see ::nvrtcAddNameExpression"] #[no_mangle] -pub unsafe extern "system" fn nvrtcGetLoweredName( +pub extern "system" fn nvrtcGetLoweredName( prog: nvrtcProgram, name_expression: *const ::std::os::raw::c_char, lowered_name: *mut *const ::std::os::raw::c_char, ) -> nvrtcResult { - crate::get_lowered_name(prog, name_expression, lowered_name) + crate::unsupported() +} + +extern crate libloading; +pub struct LibNvrtc { + __library: ::libloading::Library, + pub nvrtcGetErrorString: + unsafe extern "C" fn(result: nvrtcResult) -> *const ::std::os::raw::c_char, + pub nvrtcVersion: unsafe extern "C" fn( + major: *mut ::std::os::raw::c_int, + minor: *mut ::std::os::raw::c_int, + ) -> nvrtcResult, + pub nvrtcGetNumSupportedArchs: + unsafe extern "C" fn(numArchs: *mut ::std::os::raw::c_int) -> nvrtcResult, + pub nvrtcGetSupportedArchs: + unsafe extern "C" fn(supportedArchs: *mut ::std::os::raw::c_int) -> nvrtcResult, + pub nvrtcCreateProgram: unsafe extern "C" fn( + prog: *mut nvrtcProgram, + src: *const ::std::os::raw::c_char, + name: *const ::std::os::raw::c_char, + numHeaders: ::std::os::raw::c_int, + headers: *const *const ::std::os::raw::c_char, + includeNames: *const *const ::std::os::raw::c_char, + ) -> nvrtcResult, + pub nvrtcDestroyProgram: unsafe extern "C" fn(prog: *mut nvrtcProgram) -> nvrtcResult, + pub nvrtcCompileProgram: unsafe extern "C" fn( + prog: nvrtcProgram, + numOptions: ::std::os::raw::c_int, + options: *const *const ::std::os::raw::c_char, + ) -> nvrtcResult, + pub nvrtcGetPTXSize: + unsafe extern "C" fn(prog: nvrtcProgram, ptxSizeRet: *mut usize) -> nvrtcResult, + pub nvrtcGetPTX: + unsafe extern "C" fn(prog: nvrtcProgram, ptx: *mut ::std::os::raw::c_char) -> nvrtcResult, + pub nvrtcGetCUBINSize: + unsafe extern "C" fn(prog: nvrtcProgram, cubinSizeRet: *mut usize) -> nvrtcResult, + pub nvrtcGetCUBIN: + unsafe extern "C" fn(prog: nvrtcProgram, cubin: *mut ::std::os::raw::c_char) -> nvrtcResult, + pub nvrtcGetNVVMSize: + unsafe extern "C" fn(prog: nvrtcProgram, nvvmSizeRet: *mut usize) -> nvrtcResult, + pub nvrtcGetNVVM: + unsafe extern "C" fn(prog: nvrtcProgram, nvvm: *mut ::std::os::raw::c_char) -> nvrtcResult, + pub nvrtcGetProgramLogSize: + unsafe extern "C" fn(prog: nvrtcProgram, logSizeRet: *mut usize) -> nvrtcResult, + pub nvrtcGetProgramLog: + unsafe extern "C" fn(prog: nvrtcProgram, log: *mut ::std::os::raw::c_char) -> nvrtcResult, + pub nvrtcAddNameExpression: unsafe extern "C" fn( + prog: nvrtcProgram, + name_expression: *const ::std::os::raw::c_char, + ) -> nvrtcResult, + pub nvrtcGetLoweredName: unsafe extern "C" fn( + prog: nvrtcProgram, + name_expression: *const ::std::os::raw::c_char, + lowered_name: *mut *const ::std::os::raw::c_char, + ) -> nvrtcResult, +} +impl LibNvrtc { + pub unsafe fn new

(path: P) -> Result + where + P: AsRef<::std::ffi::OsStr>, + { + let library = ::libloading::Library::new(path)?; + Self::from_library(library) + } + pub unsafe fn from_library(library: L) -> Result + where + L: Into<::libloading::Library>, + { + let __library = library.into(); + let nvrtcGetErrorString = __library.get(b"nvrtcGetErrorString\0").map(|sym| *sym)?; + let nvrtcVersion = __library.get(b"nvrtcVersion\0").map(|sym| *sym)?; + let nvrtcGetNumSupportedArchs = __library + .get(b"nvrtcGetNumSupportedArchs\0") + .map(|sym| *sym)?; + let nvrtcGetSupportedArchs = __library.get(b"nvrtcGetSupportedArchs\0").map(|sym| *sym)?; + let nvrtcCreateProgram = __library.get(b"nvrtcCreateProgram\0").map(|sym| *sym)?; + let nvrtcDestroyProgram = __library.get(b"nvrtcDestroyProgram\0").map(|sym| *sym)?; + let nvrtcCompileProgram = __library.get(b"nvrtcCompileProgram\0").map(|sym| *sym)?; + let nvrtcGetPTXSize = __library.get(b"nvrtcGetPTXSize\0").map(|sym| *sym)?; + let nvrtcGetPTX = __library.get(b"nvrtcGetPTX\0").map(|sym| *sym)?; + let nvrtcGetCUBINSize = __library.get(b"nvrtcGetCUBINSize\0").map(|sym| *sym)?; + let nvrtcGetCUBIN = __library.get(b"nvrtcGetCUBIN\0").map(|sym| *sym)?; + let nvrtcGetNVVMSize = __library.get(b"nvrtcGetNVVMSize\0").map(|sym| *sym)?; + let nvrtcGetNVVM = __library.get(b"nvrtcGetNVVM\0").map(|sym| *sym)?; + let nvrtcGetProgramLogSize = __library.get(b"nvrtcGetProgramLogSize\0").map(|sym| *sym)?; + let nvrtcGetProgramLog = __library.get(b"nvrtcGetProgramLog\0").map(|sym| *sym)?; + let nvrtcAddNameExpression = __library.get(b"nvrtcAddNameExpression\0").map(|sym| *sym)?; + let nvrtcGetLoweredName = __library.get(b"nvrtcGetLoweredName\0").map(|sym| *sym)?; + Ok(LibNvrtc { + __library, + nvrtcGetErrorString, + nvrtcVersion, + nvrtcGetNumSupportedArchs, + nvrtcGetSupportedArchs, + nvrtcCreateProgram, + nvrtcDestroyProgram, + nvrtcCompileProgram, + nvrtcGetPTXSize, + nvrtcGetPTX, + nvrtcGetCUBINSize, + nvrtcGetCUBIN, + nvrtcGetNVVMSize, + nvrtcGetNVVM, + nvrtcGetProgramLogSize, + nvrtcGetProgramLog, + nvrtcAddNameExpression, + nvrtcGetLoweredName, + }) + } + #[doc = " \\ingroup error\n \\brief nvrtcGetErrorString is a helper function that returns a string\n describing the given nvrtcResult code, e.g., NVRTC_SUCCESS to\n \\c \"NVRTC_SUCCESS\".\n For unrecognized enumeration values, it returns\n \\c \"NVRTC_ERROR unknown\".\n\n \\param [in] result CUDA Runtime Compilation API result code.\n \\return Message string for the given #nvrtcResult code."] + pub unsafe fn nvrtcGetErrorString(&self, result: nvrtcResult) -> *const ::std::os::raw::c_char { + (self.nvrtcGetErrorString)(result) + } + #[doc = " \\ingroup query\n \\brief nvrtcVersion sets the output parameters \\p major and \\p minor\n with the CUDA Runtime Compilation version number.\n\n \\param [out] major CUDA Runtime Compilation major version number.\n \\param [out] minor CUDA Runtime Compilation minor version number.\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_INPUT \\endlink\n"] + pub unsafe fn nvrtcVersion( + &self, + major: *mut ::std::os::raw::c_int, + minor: *mut ::std::os::raw::c_int, + ) -> nvrtcResult { + (self.nvrtcVersion)(major, minor) + } + #[doc = " \\ingroup query\n \\brief nvrtcGetNumSupportedArchs sets the output parameter \\p numArchs\n with the number of architectures supported by NVRTC. This can\n then be used to pass an array to ::nvrtcGetSupportedArchs to\n get the supported architectures.\n\n \\param [out] numArchs number of supported architectures.\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_INPUT \\endlink\n\n see ::nvrtcGetSupportedArchs"] + pub unsafe fn nvrtcGetNumSupportedArchs( + &self, + numArchs: *mut ::std::os::raw::c_int, + ) -> nvrtcResult { + (self.nvrtcGetNumSupportedArchs)(numArchs) + } + #[doc = " \\ingroup query\n \\brief nvrtcGetSupportedArchs populates the array passed via the output parameter\n \\p supportedArchs with the architectures supported by NVRTC. The array is\n sorted in the ascending order. The size of the array to be passed can be\n determined using ::nvrtcGetNumSupportedArchs.\n\n \\param [out] supportedArchs sorted array of supported architectures.\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_INPUT \\endlink\n\n see ::nvrtcGetNumSupportedArchs"] + pub unsafe fn nvrtcGetSupportedArchs( + &self, + supportedArchs: *mut ::std::os::raw::c_int, + ) -> nvrtcResult { + (self.nvrtcGetSupportedArchs)(supportedArchs) + } + #[doc = " \\ingroup compilation\n \\brief nvrtcCreateProgram creates an instance of nvrtcProgram with the\n given input parameters, and sets the output parameter \\p prog with\n it.\n\n \\param [out] prog CUDA Runtime Compilation program.\n \\param [in] src CUDA program source.\n \\param [in] name CUDA program name.\\n\n \\p name can be \\c NULL; \\c \"default_program\" is\n used when \\p name is \\c NULL or \"\".\n \\param [in] numHeaders Number of headers used.\\n\n \\p numHeaders must be greater than or equal to 0.\n \\param [in] headers Sources of the headers.\\n\n \\p headers can be \\c NULL when \\p numHeaders is\n 0.\n \\param [in] includeNames Name of each header by which they can be\n included in the CUDA program source.\\n\n \\p includeNames can be \\c NULL when \\p numHeaders\n is 0.\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_OUT_OF_MEMORY \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_PROGRAM_CREATION_FAILURE \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_INPUT \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_PROGRAM \\endlink\n\n \\see ::nvrtcDestroyProgram"] + pub unsafe fn nvrtcCreateProgram( + &self, + prog: *mut nvrtcProgram, + src: *const ::std::os::raw::c_char, + name: *const ::std::os::raw::c_char, + numHeaders: ::std::os::raw::c_int, + headers: *const *const ::std::os::raw::c_char, + includeNames: *const *const ::std::os::raw::c_char, + ) -> nvrtcResult { + (self.nvrtcCreateProgram)(prog, src, name, numHeaders, headers, includeNames) + } + #[doc = " \\ingroup compilation\n \\brief nvrtcDestroyProgram destroys the given program.\n\n \\param [in] prog CUDA Runtime Compilation program.\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_PROGRAM \\endlink\n\n \\see ::nvrtcCreateProgram"] + pub unsafe fn nvrtcDestroyProgram(&self, prog: *mut nvrtcProgram) -> nvrtcResult { + (self.nvrtcDestroyProgram)(prog) + } + #[doc = " \\ingroup compilation\n \\brief nvrtcCompileProgram compiles the given program.\n\n \\param [in] prog CUDA Runtime Compilation program.\n \\param [in] numOptions Number of compiler options passed.\n \\param [in] options Compiler options in the form of C string array.\\n\n \\p options can be \\c NULL when \\p numOptions is 0.\n\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_OUT_OF_MEMORY \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_INPUT \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_PROGRAM \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_OPTION \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_COMPILATION \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_BUILTIN_OPERATION_FAILURE \\endlink\n\n It supports compile options listed in \\ref options."] + pub unsafe fn nvrtcCompileProgram( + &self, + prog: nvrtcProgram, + numOptions: ::std::os::raw::c_int, + options: *const *const ::std::os::raw::c_char, + ) -> nvrtcResult { + (self.nvrtcCompileProgram)(prog, numOptions, options) + } + #[doc = " \\ingroup compilation\n \\brief nvrtcGetPTXSize sets \\p ptxSizeRet with the size of the PTX\n generated by the previous compilation of \\p prog (including the\n trailing \\c NULL).\n\n \\param [in] prog CUDA Runtime Compilation program.\n \\param [out] ptxSizeRet Size of the generated PTX (including the trailing\n \\c NULL).\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_INPUT \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_PROGRAM \\endlink\n\n \\see ::nvrtcGetPTX"] + pub unsafe fn nvrtcGetPTXSize( + &self, + prog: nvrtcProgram, + ptxSizeRet: *mut usize, + ) -> nvrtcResult { + (self.nvrtcGetPTXSize)(prog, ptxSizeRet) + } + #[doc = " \\ingroup compilation\n \\brief nvrtcGetPTX stores the PTX generated by the previous compilation\n of \\p prog in the memory pointed by \\p ptx.\n\n \\param [in] prog CUDA Runtime Compilation program.\n \\param [out] ptx Compiled result.\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_INPUT \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_PROGRAM \\endlink\n\n \\see ::nvrtcGetPTXSize"] + pub unsafe fn nvrtcGetPTX( + &self, + prog: nvrtcProgram, + ptx: *mut ::std::os::raw::c_char, + ) -> nvrtcResult { + (self.nvrtcGetPTX)(prog, ptx) + } + #[doc = " \\ingroup compilation\n \\brief nvrtcGetCUBINSize sets \\p cubinSizeRet with the size of the cubin\n generated by the previous compilation of \\p prog. The value of\n cubinSizeRet is set to 0 if the value specified to \\c -arch is a\n virtual architecture instead of an actual architecture.\n\n \\param [in] prog CUDA Runtime Compilation program.\n \\param [out] cubinSizeRet Size of the generated cubin.\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_INPUT \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_PROGRAM \\endlink\n\n \\see ::nvrtcGetCUBIN"] + pub unsafe fn nvrtcGetCUBINSize( + &self, + prog: nvrtcProgram, + cubinSizeRet: *mut usize, + ) -> nvrtcResult { + (self.nvrtcGetCUBINSize)(prog, cubinSizeRet) + } + #[doc = " \\ingroup compilation\n \\brief nvrtcGetCUBIN stores the cubin generated by the previous compilation\n of \\p prog in the memory pointed by \\p cubin. No cubin is available\n if the value specified to \\c -arch is a virtual architecture instead\n of an actual architecture.\n\n \\param [in] prog CUDA Runtime Compilation program.\n \\param [out] cubin Compiled and assembled result.\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_INPUT \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_PROGRAM \\endlink\n\n \\see ::nvrtcGetCUBINSize"] + pub unsafe fn nvrtcGetCUBIN( + &self, + prog: nvrtcProgram, + cubin: *mut ::std::os::raw::c_char, + ) -> nvrtcResult { + (self.nvrtcGetCUBIN)(prog, cubin) + } + #[doc = " \\ingroup compilation\n \\brief nvrtcGetNVVMSize sets \\p nvvmSizeRet with the size of the NVVM\n generated by the previous compilation of \\p prog. The value of\n nvvmSizeRet is set to 0 if the program was not compiled with\n \\c -dlto.\n\n \\param [in] prog CUDA Runtime Compilation program.\n \\param [out] nvvmSizeRet Size of the generated NVVM.\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_INPUT \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_PROGRAM \\endlink\n\n \\see ::nvrtcGetNVVM"] + pub unsafe fn nvrtcGetNVVMSize( + &self, + prog: nvrtcProgram, + nvvmSizeRet: *mut usize, + ) -> nvrtcResult { + (self.nvrtcGetNVVMSize)(prog, nvvmSizeRet) + } + #[doc = " \\ingroup compilation\n \\brief nvrtcGetNVVM stores the NVVM generated by the previous compilation\n of \\p prog in the memory pointed by \\p nvvm.\n The program must have been compiled with -dlto,\n otherwise will return an error.\n\n \\param [in] prog CUDA Runtime Compilation program.\n \\param [out] nvvm Compiled result.\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_INPUT \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_PROGRAM \\endlink\n\n \\see ::nvrtcGetNVVMSize"] + pub unsafe fn nvrtcGetNVVM( + &self, + prog: nvrtcProgram, + nvvm: *mut ::std::os::raw::c_char, + ) -> nvrtcResult { + (self.nvrtcGetNVVM)(prog, nvvm) + } + #[doc = " \\ingroup compilation\n \\brief nvrtcGetProgramLogSize sets \\p logSizeRet with the size of the\n log generated by the previous compilation of \\p prog (including the\n trailing \\c NULL).\n\n Note that compilation log may be generated with warnings and informative\n messages, even when the compilation of \\p prog succeeds.\n\n \\param [in] prog CUDA Runtime Compilation program.\n \\param [out] logSizeRet Size of the compilation log\n (including the trailing \\c NULL).\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_INPUT \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_PROGRAM \\endlink\n\n \\see ::nvrtcGetProgramLog"] + pub unsafe fn nvrtcGetProgramLogSize( + &self, + prog: nvrtcProgram, + logSizeRet: *mut usize, + ) -> nvrtcResult { + (self.nvrtcGetProgramLogSize)(prog, logSizeRet) + } + #[doc = " \\ingroup compilation\n \\brief nvrtcGetProgramLog stores the log generated by the previous\n compilation of \\p prog in the memory pointed by \\p log.\n\n \\param [in] prog CUDA Runtime Compilation program.\n \\param [out] log Compilation log.\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_INPUT \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_INVALID_PROGRAM \\endlink\n\n \\see ::nvrtcGetProgramLogSize"] + pub unsafe fn nvrtcGetProgramLog( + &self, + prog: nvrtcProgram, + log: *mut ::std::os::raw::c_char, + ) -> nvrtcResult { + (self.nvrtcGetProgramLog)(prog, log) + } + #[doc = " \\ingroup compilation\n \\brief nvrtcAddNameExpression notes the given name expression\n denoting the address of a __global__ function\n or __device__/__constant__ variable.\n\n The identical name expression string must be provided on a subsequent\n call to nvrtcGetLoweredName to extract the lowered name.\n \\param [in] prog CUDA Runtime Compilation program.\n \\param [in] name_expression constant expression denoting the address of\n a __global__ function or __device__/__constant__ variable.\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION \\endlink\n\n \\see ::nvrtcGetLoweredName"] + pub unsafe fn nvrtcAddNameExpression( + &self, + prog: nvrtcProgram, + name_expression: *const ::std::os::raw::c_char, + ) -> nvrtcResult { + (self.nvrtcAddNameExpression)(prog, name_expression) + } + #[doc = " \\ingroup compilation\n \\brief nvrtcGetLoweredName extracts the lowered (mangled) name\n for a __global__ function or __device__/__constant__ variable,\n and updates *lowered_name to point to it. The memory containing\n the name is released when the NVRTC program is destroyed by\n nvrtcDestroyProgram.\n The identical name expression must have been previously\n provided to nvrtcAddNameExpression.\n\n \\param [in] prog CUDA Runtime Compilation program.\n \\param [in] name_expression constant expression denoting the address of\n a __global__ function or __device__/__constant__ variable.\n \\param [out] lowered_name initialized by the function to point to a\n C string containing the lowered (mangled)\n name corresponding to the provided name expression.\n \\return\n - \\link #nvrtcResult NVRTC_SUCCESS \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION \\endlink\n - \\link #nvrtcResult NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID \\endlink\n\n \\see ::nvrtcAddNameExpression"] + pub unsafe fn nvrtcGetLoweredName( + &self, + prog: nvrtcProgram, + name_expression: *const ::std::os::raw::c_char, + lowered_name: *mut *const ::std::os::raw::c_char, + ) -> nvrtcResult { + (self.nvrtcGetLoweredName)(prog, name_expression, lowered_name) + } }