Implement runtime compiler. (nvrtc)
This commit is contained in:
Seunghoon Lee
parent
3adbe3ec30
commit
7d33957a85
|
|
@ -14,6 +14,7 @@ members = [
|
|||
"hipblaslt-sys",
|
||||
"hipfft-sys",
|
||||
"hiprt-sys",
|
||||
"hiprtc-sys",
|
||||
"miopen-sys",
|
||||
"offline_compiler",
|
||||
"optix_base",
|
||||
|
|
@ -39,6 +40,7 @@ members = [
|
|||
"zluda_ml",
|
||||
"zluda_redirect",
|
||||
"zluda_rt",
|
||||
"zluda_rtc",
|
||||
"zluda_sparse",
|
||||
]
|
||||
|
||||
|
|
|
|||
|
|
@ -22,6 +22,7 @@ args = [
|
|||
"-p", "zluda_lib",
|
||||
"-p", "zluda_ml",
|
||||
"-p", "zluda_sparse",
|
||||
"-p", "zluda_rtc",
|
||||
"-p", "zluda_redirect",
|
||||
]
|
||||
|
||||
|
|
@ -38,6 +39,7 @@ args = [
|
|||
"-p", "zluda_fft",
|
||||
"-p", "zluda_lib",
|
||||
"-p", "zluda_ml",
|
||||
"-p", "zluda_rtc",
|
||||
"-p", "zluda_sparse",
|
||||
]
|
||||
|
||||
|
|
@ -55,6 +57,7 @@ args = [
|
|||
"-p", "zluda_fft",
|
||||
"-p", "zluda_lib",
|
||||
"-p", "zluda_ml",
|
||||
"-p", "zluda_rtc",
|
||||
"-p", "zluda_sparse",
|
||||
]
|
||||
|
||||
|
|
|
|||
8
hiprtc-sys/Cargo.toml
Normal file
8
hiprtc-sys/Cargo.toml
Normal file
|
|
@ -0,0 +1,8 @@
|
|||
[package]
|
||||
name = "hiprtc-sys"
|
||||
version = "0.0.0"
|
||||
authors = ["Seunghoon Lee <op@lsh.sh>"]
|
||||
edition = "2018"
|
||||
links = "hiprtc"
|
||||
|
||||
[lib]
|
||||
1
hiprtc-sys/README
Normal file
1
hiprtc-sys/README
Normal file
|
|
@ -0,0 +1 @@
|
|||
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__
|
||||
14
hiprtc-sys/build.rs
Normal file
14
hiprtc-sys/build.rs
Normal file
|
|
@ -0,0 +1,14 @@
|
|||
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(())
|
||||
}
|
||||
360
hiprtc-sys/src/hiprtc.rs
Normal file
360
hiprtc-sys/src/hiprtc.rs
Normal file
|
|
@ -0,0 +1,360 @@
|
|||
/* 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"]
|
||||
pub const HIPRTC_JIT_MAX_REGISTERS: hiprtcJIT_option = hiprtcJIT_option(0);
|
||||
}
|
||||
impl hiprtcJIT_option {
|
||||
#[doc = "< Thread per block"]
|
||||
pub const HIPRTC_JIT_THREADS_PER_BLOCK: hiprtcJIT_option = hiprtcJIT_option(1);
|
||||
}
|
||||
impl hiprtcJIT_option {
|
||||
#[doc = "< Time from aall clock"]
|
||||
pub const HIPRTC_JIT_WALL_TIME: hiprtcJIT_option = hiprtcJIT_option(2);
|
||||
}
|
||||
impl hiprtcJIT_option {
|
||||
#[doc = "< Log buffer info"]
|
||||
pub const HIPRTC_JIT_INFO_LOG_BUFFER: hiprtcJIT_option = hiprtcJIT_option(3);
|
||||
}
|
||||
impl hiprtcJIT_option {
|
||||
#[doc = "< Log buffer size in bytes"]
|
||||
pub const HIPRTC_JIT_INFO_LOG_BUFFER_SIZE_BYTES: hiprtcJIT_option = hiprtcJIT_option(4);
|
||||
}
|
||||
impl hiprtcJIT_option {
|
||||
#[doc = "< Log buffer error"]
|
||||
pub const HIPRTC_JIT_ERROR_LOG_BUFFER: hiprtcJIT_option = hiprtcJIT_option(5);
|
||||
}
|
||||
impl hiprtcJIT_option {
|
||||
#[doc = "< Log buffer size in bytes"]
|
||||
pub const HIPRTC_JIT_ERROR_LOG_BUFFER_SIZE_BYTES: hiprtcJIT_option = hiprtcJIT_option(6);
|
||||
}
|
||||
impl hiprtcJIT_option {
|
||||
#[doc = "< Optimization level"]
|
||||
pub const HIPRTC_JIT_OPTIMIZATION_LEVEL: hiprtcJIT_option = hiprtcJIT_option(7);
|
||||
}
|
||||
impl hiprtcJIT_option {
|
||||
#[doc = "<"]
|
||||
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 = "< Log verbose"]
|
||||
pub const HIPRTC_JIT_LOG_VERBOSE: hiprtcJIT_option = hiprtcJIT_option(12);
|
||||
}
|
||||
impl hiprtcJIT_option {
|
||||
#[doc = "< Generate line information"]
|
||||
pub const HIPRTC_JIT_GENERATE_LINE_INFO: hiprtcJIT_option = hiprtcJIT_option(13);
|
||||
}
|
||||
impl hiprtcJIT_option {
|
||||
#[doc = "< Cache mode"]
|
||||
pub const HIPRTC_JIT_CACHE_MODE: hiprtcJIT_option = hiprtcJIT_option(14);
|
||||
}
|
||||
impl hiprtcJIT_option {
|
||||
#[doc = "< New SM3X option"]
|
||||
pub const HIPRTC_JIT_NEW_SM3X_OPT: hiprtcJIT_option = hiprtcJIT_option(15);
|
||||
}
|
||||
impl hiprtcJIT_option {
|
||||
#[doc = "< Fast compile"]
|
||||
pub const HIPRTC_JIT_FAST_COMPILE: hiprtcJIT_option = hiprtcJIT_option(16);
|
||||
}
|
||||
impl hiprtcJIT_option {
|
||||
#[doc = "< Global symbol names"]
|
||||
pub const HIPRTC_JIT_GLOBAL_SYMBOL_NAMES: hiprtcJIT_option = hiprtcJIT_option(17);
|
||||
}
|
||||
impl hiprtcJIT_option {
|
||||
#[doc = "< Global symbol address"]
|
||||
pub const HIPRTC_JIT_GLOBAL_SYMBOL_ADDRESS: hiprtcJIT_option = hiprtcJIT_option(18);
|
||||
}
|
||||
impl hiprtcJIT_option {
|
||||
#[doc = "< Global symbol count"]
|
||||
pub const HIPRTC_JIT_GLOBAL_SYMBOL_COUNT: hiprtcJIT_option = hiprtcJIT_option(19);
|
||||
}
|
||||
impl hiprtcJIT_option {
|
||||
#[doc = "< LTO"]
|
||||
pub const HIPRTC_JIT_LTO: hiprtcJIT_option = hiprtcJIT_option(20);
|
||||
}
|
||||
impl hiprtcJIT_option {
|
||||
#[doc = "< FTZ"]
|
||||
pub const HIPRTC_JIT_FTZ: hiprtcJIT_option = hiprtcJIT_option(21);
|
||||
}
|
||||
impl hiprtcJIT_option {
|
||||
#[doc = "< Prec_VIV"]
|
||||
pub const HIPRTC_JIT_PREC_DIV: hiprtcJIT_option = hiprtcJIT_option(22);
|
||||
}
|
||||
impl hiprtcJIT_option {
|
||||
#[doc = "< PREC_SQRT"]
|
||||
pub const HIPRTC_JIT_PREC_SQRT: hiprtcJIT_option = hiprtcJIT_option(23);
|
||||
}
|
||||
impl hiprtcJIT_option {
|
||||
#[doc = "< FMA"]
|
||||
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 = "< AMD only. Linker options to be passed on to"]
|
||||
pub const HIPRTC_JIT_IR_TO_ISA_OPT_EXT: hiprtcJIT_option = hiprtcJIT_option(10000);
|
||||
}
|
||||
impl hiprtcJIT_option {
|
||||
#[doc = "< AMD only. Count of linker options"]
|
||||
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"]
|
||||
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;
|
||||
}
|
||||
3
hiprtc-sys/src/lib.rs
Normal file
3
hiprtc-sys/src/lib.rs
Normal file
|
|
@ -0,0 +1,3 @@
|
|||
#[allow(warnings)]
|
||||
mod hiprtc;
|
||||
pub use hiprtc::*;
|
||||
|
|
@ -23,11 +23,8 @@ use winapi::um::{
|
|||
use winapi::um::winbase::{INFINITE, WAIT_FAILED};
|
||||
|
||||
static REDIRECT_DLL: &'static str = "zluda_redirect.dll";
|
||||
static CUBLAS_DLL: &'static str = "cublas.dll";
|
||||
static CUDNN_DLL: &'static str = "cudnn.dll";
|
||||
static CUFFT_DLL: &'static str = "cufft.dll";
|
||||
static CUSPARSE_DLL: &'static str = "cusparse.dll";
|
||||
static NCCL_DLL: &'static str = "nccl.dll";
|
||||
static NVRTC_DLL: &'static str = "nvrtc.dll";
|
||||
static NVCUDA_DLL: &'static str = "nvcuda.dll";
|
||||
static NVML_DLL: &'static str = "nvml.dll";
|
||||
static NVAPI_DLL: &'static str = "nvapi64.dll";
|
||||
|
|
@ -38,26 +35,14 @@ include!("../../zluda_redirect/src/payload_guid.rs");
|
|||
#[derive(FromArgs)]
|
||||
/// Launch application with custom CUDA libraries
|
||||
struct ProgramArguments {
|
||||
/// DLL to be injected instead of system cublas.dll. If not provided {0}, will use cublas.dll from its own directory
|
||||
#[argh(option)]
|
||||
cublas: Option<PathBuf>,
|
||||
|
||||
/// DLL to be injected instead of system cudnn.dll. If not provided {0}, will use cudnn.dll from its own directory
|
||||
#[argh(option)]
|
||||
cudnn: Option<PathBuf>,
|
||||
|
||||
/// DLL to be injected instead of system cufft.dll. If not provided {0}, will use cufft.dll from its own directory
|
||||
#[argh(option)]
|
||||
cufft: Option<PathBuf>,
|
||||
|
||||
/// DLL to be injected instead of system cusparse.dll. If not provided {0}, will use cusparse.dll from its own directory
|
||||
#[argh(option)]
|
||||
cusparse: Option<PathBuf>,
|
||||
|
||||
/// DLL to be injected instead of system nccl.dll. If not provided {0}, will use nccl.dll from its own directory
|
||||
#[argh(option)]
|
||||
nccl: Option<PathBuf>,
|
||||
|
||||
/// DLL to be injected instead of system nvrtc.dll. If not provided {0}, will use nvrtc.dll from its own directory
|
||||
#[argh(option)]
|
||||
nvrtc: Option<PathBuf>,
|
||||
|
||||
/// DLL to be injected instead of system nvcuda.dll. If not provided {0}, will use nvcuda.dll from its own directory
|
||||
#[argh(option)]
|
||||
nvcuda: Option<PathBuf>,
|
||||
|
|
@ -90,11 +75,8 @@ pub fn main_impl() -> Result<(), Box<dyn Error>> {
|
|||
let mut startup_info = unsafe { mem::zeroed::<detours_sys::_STARTUPINFOW>() };
|
||||
let mut proc_info = unsafe { mem::zeroed::<detours_sys::_PROCESS_INFORMATION>() };
|
||||
let mut dlls_to_inject = vec![
|
||||
environment.cublas_path_zero_terminated.as_ptr() as _,
|
||||
//environment.cudnn_path_zero_terminated.as_ptr() as _,
|
||||
environment.cufft_path_zero_terminated.as_ptr() as _,
|
||||
environment.cusparse_path_zero_terminated.as_ptr() as _,
|
||||
environment.nccl_path_zero_terminated.as_ptr() as _,
|
||||
environment.nvrtc_path_zero_terminated.as_ptr() as _,
|
||||
environment.nvcuda_path_zero_terminated.as_ptr() as _,
|
||||
environment.nvml_path_zero_terminated.as_ptr() as *const i8,
|
||||
environment.redirect_path_zero_terminated.as_ptr() as _,
|
||||
|
|
@ -176,11 +158,8 @@ pub fn main_impl() -> Result<(), Box<dyn Error>> {
|
|||
}
|
||||
|
||||
struct NormalizedArguments {
|
||||
cublas_path: PathBuf,
|
||||
cudnn_path: PathBuf,
|
||||
cufft_path: PathBuf,
|
||||
cusparse_path: PathBuf,
|
||||
nccl_path: PathBuf,
|
||||
nvrtc_path: PathBuf,
|
||||
nvcuda_path: PathBuf,
|
||||
nvml_path: PathBuf,
|
||||
nvapi_path: Option<PathBuf>,
|
||||
|
|
@ -192,16 +171,10 @@ struct NormalizedArguments {
|
|||
impl NormalizedArguments {
|
||||
fn new(prog_args: ProgramArguments) -> Result<Self, Box<dyn Error>> {
|
||||
let current_exe = env::current_exe()?;
|
||||
let cublas_path =
|
||||
Self::get_absolute_path_or_default(¤t_exe, prog_args.cublas, CUBLAS_DLL)?;
|
||||
let cudnn_path =
|
||||
Self::get_absolute_path_or_default(¤t_exe, prog_args.cudnn, CUDNN_DLL)?;
|
||||
let cufft_path =
|
||||
Self::get_absolute_path_or_default(¤t_exe, prog_args.cufft, CUFFT_DLL)?;
|
||||
let cusparse_path =
|
||||
Self::get_absolute_path_or_default(¤t_exe, prog_args.cusparse, CUSPARSE_DLL)?;
|
||||
let nccl_path =
|
||||
Self::get_absolute_path_or_default(¤t_exe, prog_args.nccl, NCCL_DLL)?;
|
||||
let nvrtc_path =
|
||||
Self::get_absolute_path_or_default(¤t_exe, prog_args.nvrtc, NVRTC_DLL)?;
|
||||
let nvcuda_path =
|
||||
Self::get_absolute_path_or_default(¤t_exe, prog_args.nvcuda, NVCUDA_DLL)?;
|
||||
let nvml_path = Self::get_absolute_path_or_default(¤t_exe, prog_args.nvml, NVML_DLL)?;
|
||||
|
|
@ -212,11 +185,8 @@ impl NormalizedArguments {
|
|||
let mut redirect_path = current_exe.parent().unwrap().to_path_buf();
|
||||
redirect_path.push(REDIRECT_DLL);
|
||||
Ok(Self {
|
||||
cublas_path,
|
||||
cudnn_path,
|
||||
cufft_path,
|
||||
cusparse_path,
|
||||
nccl_path,
|
||||
nvrtc_path,
|
||||
nvcuda_path,
|
||||
nvml_path,
|
||||
nvapi_path,
|
||||
|
|
@ -274,11 +244,8 @@ impl NormalizedArguments {
|
|||
}
|
||||
|
||||
struct Environment {
|
||||
cublas_path_zero_terminated: String,
|
||||
cudnn_path_zero_terminated: String,
|
||||
cufft_path_zero_terminated: String,
|
||||
cusparse_path_zero_terminated: String,
|
||||
nccl_path_zero_terminated: String,
|
||||
nvrtc_path_zero_terminated: String,
|
||||
nvcuda_path_zero_terminated: String,
|
||||
nvml_path_zero_terminated: String,
|
||||
nvapi_path_zero_terminated: Option<String>,
|
||||
|
|
@ -294,31 +261,16 @@ struct Environment {
|
|||
impl Environment {
|
||||
fn setup(args: NormalizedArguments) -> io::Result<Self> {
|
||||
let _temp_dir = TempDir::new()?;
|
||||
let cublas_path_zero_terminated = Self::zero_terminate(Self::copy_to_correct_name(
|
||||
args.cublas_path,
|
||||
&_temp_dir,
|
||||
CUBLAS_DLL,
|
||||
)?);
|
||||
let cudnn_path_zero_terminated = Self::zero_terminate(Self::copy_to_correct_name(
|
||||
args.cudnn_path,
|
||||
&_temp_dir,
|
||||
CUDNN_DLL,
|
||||
)?);
|
||||
let cufft_path_zero_terminated = Self::zero_terminate(Self::copy_to_correct_name(
|
||||
args.cufft_path,
|
||||
&_temp_dir,
|
||||
CUFFT_DLL,
|
||||
)?);
|
||||
let cusparse_path_zero_terminated = Self::zero_terminate(Self::copy_to_correct_name(
|
||||
args.cusparse_path,
|
||||
&_temp_dir,
|
||||
CUSPARSE_DLL,
|
||||
)?);
|
||||
let nccl_path_zero_terminated = Self::zero_terminate(Self::copy_to_correct_name(
|
||||
args.nccl_path,
|
||||
&_temp_dir,
|
||||
NCCL_DLL,
|
||||
)?);
|
||||
let nvrtc_path_zero_terminated = Self::zero_terminate(Self::copy_to_correct_name(
|
||||
args.nvrtc_path,
|
||||
&_temp_dir,
|
||||
NVRTC_DLL,
|
||||
)?);
|
||||
let nvcuda_path_zero_terminated = Self::zero_terminate(Self::copy_to_correct_name(
|
||||
args.nvcuda_path,
|
||||
&_temp_dir,
|
||||
|
|
@ -349,11 +301,8 @@ impl Environment {
|
|||
.transpose()?;
|
||||
let redirect_path_zero_terminated = Self::zero_terminate(args.redirect_path);
|
||||
Ok(Self {
|
||||
cublas_path_zero_terminated,
|
||||
cudnn_path_zero_terminated,
|
||||
cufft_path_zero_terminated,
|
||||
cusparse_path_zero_terminated,
|
||||
nccl_path_zero_terminated,
|
||||
nvrtc_path_zero_terminated,
|
||||
nvcuda_path_zero_terminated,
|
||||
nvml_path_zero_terminated,
|
||||
nvapi_path_zero_terminated,
|
||||
|
|
|
|||
|
|
@ -52,10 +52,6 @@ use winapi::{
|
|||
include!("payload_guid.rs");
|
||||
|
||||
const WIN_MAX_PATH: usize = 260;
|
||||
const CUBLAS_UTF8: &'static str = "CUBLAS.DLL";
|
||||
const CUBLAS_UTF16: &[u16] = wch!("CUBLAS.DLL");
|
||||
const CUDNN_UTF8: &'static str = "CUDNN.DLL";
|
||||
const CUDNN_UTF16: &[u16] = wch!("CUDNN.DLL");
|
||||
const NVCUDA1_UTF8: &'static str = "NVCUDA.DLL";
|
||||
const NVCUDA1_UTF16: &[u16] = wch!("NVCUDA.DLL");
|
||||
const NVCUDA2_UTF8: &'static str = "NVCUDA.DLL";
|
||||
|
|
@ -68,10 +64,6 @@ const NVOPTIX_UTF8: &'static str = "OPTIX.6.6.0.DLL";
|
|||
const NVOPTIX_UTF16: &[u16] = wch!("OPTIX.6.6.0.DLL");
|
||||
static mut ZLUDA_PATH_UTF8: Option<&'static [u8]> = None;
|
||||
static mut ZLUDA_PATH_UTF16: Vec<u16> = Vec::new();
|
||||
static mut ZLUDA_BLAS_PATH_UTF8: Option<&'static [u8]> = None;
|
||||
static mut ZLUDA_BLAS_PATH_UTF16: Vec<u16> = Vec::new();
|
||||
static mut ZLUDA_DNN_PATH_UTF8: Option<&'static [u8]> = None;
|
||||
static mut ZLUDA_DNN_PATH_UTF16: Vec<u16> = Vec::new();
|
||||
static mut ZLUDA_ML_PATH_UTF8: Option<&'static [u8]> = None;
|
||||
static mut ZLUDA_ML_PATH_UTF16: Vec<u16> = Vec::new();
|
||||
static mut ZLUDA_API_PATH_UTF8: Option<&'static [u8]> = None;
|
||||
|
|
@ -207,11 +199,7 @@ unsafe fn get_library_name_utf8(raw_library_name: *const u8) -> *const u8 {
|
|||
}
|
||||
}
|
||||
}
|
||||
if is_cublas_dll_utf8(library_name) {
|
||||
return ZLUDA_BLAS_PATH_UTF8.unwrap().as_ptr();
|
||||
} /*else if is_cudnn_dll_utf8(library_name) {
|
||||
return ZLUDA_DNN_PATH_UTF8.unwrap().as_ptr();
|
||||
}*/ else if is_nvcuda_dll_utf8(library_name) {
|
||||
if is_nvcuda_dll_utf8(library_name) {
|
||||
return ZLUDA_PATH_UTF8.unwrap().as_ptr();
|
||||
} else if is_nvml_dll_utf8(library_name) {
|
||||
return ZLUDA_ML_PATH_UTF8.unwrap().as_ptr();
|
||||
|
|
@ -249,11 +237,7 @@ unsafe fn get_library_name_utf16(raw_library_name: *const u16) -> *const u16 {
|
|||
}
|
||||
}
|
||||
}
|
||||
if is_cublas_dll_utf16(library_name) {
|
||||
return ZLUDA_BLAS_PATH_UTF16.as_ptr();
|
||||
} /*else if is_cudnn_dll_utf16(library_name) {
|
||||
return ZLUDA_DNN_PATH_UTF16.as_ptr();
|
||||
}*/ else if is_nvcuda_dll_utf16(library_name) {
|
||||
if is_nvcuda_dll_utf16(library_name) {
|
||||
return ZLUDA_PATH_UTF16.as_ptr();
|
||||
} else if is_nvml_dll_utf16(library_name) {
|
||||
return ZLUDA_ML_PATH_UTF16.as_ptr();
|
||||
|
|
@ -329,22 +313,6 @@ unsafe fn is_driverstore_utf16(lib: &[u16]) -> bool {
|
|||
starts_with_ignore_case(lib, &DRIVERSTORE_UTF16, utf16_to_ascii_uppercase)
|
||||
}
|
||||
|
||||
fn is_cublas_dll_utf8(lib: &[u8]) -> bool {
|
||||
is_dll_utf8(lib, CUBLAS_UTF8.as_bytes())
|
||||
}
|
||||
|
||||
fn is_cublas_dll_utf16(lib: &[u16]) -> bool {
|
||||
is_dll_utf16(lib, CUBLAS_UTF16)
|
||||
}
|
||||
|
||||
fn is_cudnn_dll_utf8(lib: &[u8]) -> bool {
|
||||
is_dll_utf8(lib, CUDNN_UTF8.as_bytes())
|
||||
}
|
||||
|
||||
fn is_cudnn_dll_utf16(lib: &[u16]) -> bool {
|
||||
is_dll_utf16(lib, CUDNN_UTF16)
|
||||
}
|
||||
|
||||
fn is_nvcuda_dll_utf8(lib: &[u8]) -> bool {
|
||||
is_dll_utf8(lib, NVCUDA1_UTF8.as_bytes()) || is_dll_utf8(lib, NVCUDA2_UTF8.as_bytes())
|
||||
}
|
||||
|
|
|
|||
17
zluda_rtc/Cargo.toml
Normal file
17
zluda_rtc/Cargo.toml
Normal file
|
|
@ -0,0 +1,17 @@
|
|||
[package]
|
||||
name = "zluda_rtc"
|
||||
version = "0.0.0"
|
||||
authors = ["Seunghoon Lee <op@lsh.sh>"]
|
||||
edition = "2018"
|
||||
|
||||
[lib]
|
||||
name = "nvrtc"
|
||||
crate-type = ["cdylib"]
|
||||
|
||||
[dependencies]
|
||||
hip_common = { path = "../hip_common" }
|
||||
hiprtc-sys = { path = "../hiprtc-sys" }
|
||||
|
||||
[package.metadata.zluda]
|
||||
linux_names = ["libnvrtc.so.10", "libnvrtc.so.11"]
|
||||
dump_names = ["libnvrtc.so"]
|
||||
3
zluda_rtc/README
Normal file
3
zluda_rtc/README
Normal file
|
|
@ -0,0 +1,3 @@
|
|||
bindgen include/nvrtc.h -o src/nvrtc.rs --allowlist-function="^nvrtc.*" --default-enum-style=newtype --no-layout-tests --no-derive-debug -- -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
|
||||
53
zluda_rtc/src/lib.rs
Normal file
53
zluda_rtc/src/lib.rs
Normal file
|
|
@ -0,0 +1,53 @@
|
|||
mod nvrtc;
|
||||
pub use nvrtc::*;
|
||||
|
||||
use hiprtc_sys::*;
|
||||
|
||||
#[cfg(debug_assertions)]
|
||||
fn unsupported() -> nvrtcResult {
|
||||
unimplemented!()
|
||||
}
|
||||
|
||||
#[cfg(not(debug_assertions))]
|
||||
fn unsupported() -> nvrtcResult {
|
||||
nvrtcResult::NVRTC_ERROR_INTERNAL_ERROR
|
||||
}
|
||||
|
||||
fn to_nvrtc(status: hiprtc_sys::hiprtcResult) -> nvrtcResult {
|
||||
match status {
|
||||
hiprtc_sys::hiprtcResult::HIPRTC_SUCCESS => nvrtcResult::NVRTC_SUCCESS,
|
||||
err => panic!("[ZLUDA] HIPRTC failed: {}", err.0),
|
||||
}
|
||||
}
|
||||
|
||||
unsafe fn create_program(
|
||||
prog: *mut nvrtcProgram,
|
||||
src: *const std::ffi::c_char,
|
||||
name: *const std::ffi::c_char,
|
||||
num_headers: i32,
|
||||
headers: *const *const std::ffi::c_char,
|
||||
include_names: *const *const std::ffi::c_char,
|
||||
) -> nvrtcResult {
|
||||
to_nvrtc(hiprtcCreateProgram(
|
||||
prog.cast(),
|
||||
src,
|
||||
name,
|
||||
num_headers,
|
||||
headers.cast_mut(),
|
||||
include_names.cast_mut(),
|
||||
))
|
||||
}
|
||||
|
||||
unsafe fn destroy_program(
|
||||
prog: *mut nvrtcProgram,
|
||||
) -> nvrtcResult {
|
||||
to_nvrtc(hiprtcDestroyProgram(prog.cast()))
|
||||
}
|
||||
|
||||
unsafe fn compile_program(
|
||||
prog: nvrtcProgram,
|
||||
num_options: i32,
|
||||
options: *const *const std::ffi::c_char,
|
||||
) -> nvrtcResult {
|
||||
to_nvrtc(hiprtcCompileProgram(prog.cast(), num_options, options.cast_mut()))
|
||||
}
|
||||
197
zluda_rtc/src/nvrtc.rs
Normal file
197
zluda_rtc/src/nvrtc.rs
Normal file
|
|
@ -0,0 +1,197 @@
|
|||
/* automatically generated by rust-bindgen 0.69.4 */
|
||||
|
||||
impl nvrtcResult {
|
||||
pub const NVRTC_SUCCESS: nvrtcResult = nvrtcResult(0);
|
||||
}
|
||||
impl nvrtcResult {
|
||||
pub const NVRTC_ERROR_OUT_OF_MEMORY: nvrtcResult = nvrtcResult(1);
|
||||
}
|
||||
impl nvrtcResult {
|
||||
pub const NVRTC_ERROR_PROGRAM_CREATION_FAILURE: nvrtcResult = nvrtcResult(2);
|
||||
}
|
||||
impl nvrtcResult {
|
||||
pub const NVRTC_ERROR_INVALID_INPUT: nvrtcResult = nvrtcResult(3);
|
||||
}
|
||||
impl nvrtcResult {
|
||||
pub const NVRTC_ERROR_INVALID_PROGRAM: nvrtcResult = nvrtcResult(4);
|
||||
}
|
||||
impl nvrtcResult {
|
||||
pub const NVRTC_ERROR_INVALID_OPTION: nvrtcResult = nvrtcResult(5);
|
||||
}
|
||||
impl nvrtcResult {
|
||||
pub const NVRTC_ERROR_COMPILATION: nvrtcResult = nvrtcResult(6);
|
||||
}
|
||||
impl nvrtcResult {
|
||||
pub const NVRTC_ERROR_BUILTIN_OPERATION_FAILURE: nvrtcResult = nvrtcResult(7);
|
||||
}
|
||||
impl nvrtcResult {
|
||||
pub const NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION: nvrtcResult = nvrtcResult(8);
|
||||
}
|
||||
impl nvrtcResult {
|
||||
pub const NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION: nvrtcResult = nvrtcResult(9);
|
||||
}
|
||||
impl nvrtcResult {
|
||||
pub const NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID: nvrtcResult = nvrtcResult(10);
|
||||
}
|
||||
impl nvrtcResult {
|
||||
pub const NVRTC_ERROR_INTERNAL_ERROR: nvrtcResult = nvrtcResult(11);
|
||||
}
|
||||
#[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)]
|
||||
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 extern "system" fn nvrtcGetErrorString(result: nvrtcResult) -> *const ::std::os::raw::c_char {
|
||||
unimplemented!()
|
||||
}
|
||||
|
||||
#[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"]
|
||||
#[no_mangle]
|
||||
pub extern "system" fn nvrtcVersion(
|
||||
major: *mut ::std::os::raw::c_int,
|
||||
minor: *mut ::std::os::raw::c_int,
|
||||
) -> nvrtcResult {
|
||||
crate::unsupported()
|
||||
}
|
||||
|
||||
#[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"]
|
||||
#[no_mangle]
|
||||
pub extern "system" fn nvrtcGetNumSupportedArchs(
|
||||
numArchs: *mut ::std::os::raw::c_int,
|
||||
) -> nvrtcResult {
|
||||
crate::unsupported()
|
||||
}
|
||||
|
||||
#[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"]
|
||||
#[no_mangle]
|
||||
pub extern "system" fn nvrtcGetSupportedArchs(
|
||||
supportedArchs: *mut ::std::os::raw::c_int,
|
||||
) -> nvrtcResult {
|
||||
crate::unsupported()
|
||||
}
|
||||
#[repr(C)]
|
||||
#[derive(Copy, Clone)]
|
||||
pub struct _nvrtcProgram {
|
||||
_unused: [u8; 0],
|
||||
}
|
||||
#[doc = " \\ingroup compilation\n \\brief nvrtcProgram is the unit of compilation, and an opaque handle for\n a program.\n\n To compile a CUDA program string, an instance of nvrtcProgram must be\n created first with ::nvrtcCreateProgram, then compiled with\n ::nvrtcCompileProgram."]
|
||||
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(
|
||||
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 {
|
||||
crate::create_program(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"]
|
||||
#[no_mangle]
|
||||
pub unsafe 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(
|
||||
prog: nvrtcProgram,
|
||||
numOptions: ::std::os::raw::c_int,
|
||||
options: *const *const ::std::os::raw::c_char,
|
||||
) -> nvrtcResult {
|
||||
crate::compile_program(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"]
|
||||
#[no_mangle]
|
||||
pub extern "system" fn nvrtcGetPTXSize(prog: nvrtcProgram, ptxSizeRet: *mut usize) -> nvrtcResult {
|
||||
crate::unsupported()
|
||||
}
|
||||
|
||||
#[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"]
|
||||
#[no_mangle]
|
||||
pub extern "system" fn nvrtcGetPTX(
|
||||
prog: nvrtcProgram,
|
||||
ptx: *mut ::std::os::raw::c_char,
|
||||
) -> nvrtcResult {
|
||||
crate::unsupported()
|
||||
}
|
||||
|
||||
#[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 extern "system" fn nvrtcGetCUBINSize(
|
||||
prog: nvrtcProgram,
|
||||
cubinSizeRet: *mut usize,
|
||||
) -> nvrtcResult {
|
||||
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 extern "system" fn nvrtcGetCUBIN(
|
||||
prog: nvrtcProgram,
|
||||
cubin: *mut ::std::os::raw::c_char,
|
||||
) -> nvrtcResult {
|
||||
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"]
|
||||
#[no_mangle]
|
||||
pub extern "system" fn nvrtcGetNVVMSize(
|
||||
prog: nvrtcProgram,
|
||||
nvvmSizeRet: *mut usize,
|
||||
) -> nvrtcResult {
|
||||
crate::unsupported()
|
||||
}
|
||||
|
||||
#[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"]
|
||||
#[no_mangle]
|
||||
pub extern "system" fn nvrtcGetNVVM(
|
||||
prog: nvrtcProgram,
|
||||
nvvm: *mut ::std::os::raw::c_char,
|
||||
) -> nvrtcResult {
|
||||
crate::unsupported()
|
||||
}
|
||||
|
||||
#[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"]
|
||||
#[no_mangle]
|
||||
pub extern "system" fn nvrtcGetProgramLogSize(
|
||||
prog: nvrtcProgram,
|
||||
logSizeRet: *mut usize,
|
||||
) -> nvrtcResult {
|
||||
crate::unsupported()
|
||||
}
|
||||
|
||||
#[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"]
|
||||
#[no_mangle]
|
||||
pub extern "system" fn nvrtcGetProgramLog(
|
||||
prog: nvrtcProgram,
|
||||
log: *mut ::std::os::raw::c_char,
|
||||
) -> nvrtcResult {
|
||||
crate::unsupported()
|
||||
}
|
||||
|
||||
#[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"]
|
||||
#[no_mangle]
|
||||
pub extern "system" fn nvrtcAddNameExpression(
|
||||
prog: nvrtcProgram,
|
||||
name_expression: *const ::std::os::raw::c_char,
|
||||
) -> nvrtcResult {
|
||||
crate::unsupported()
|
||||
}
|
||||
|
||||
#[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 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::unsupported()
|
||||
}
|
||||
Loading…
Reference in a new issue