diff --git a/Cargo.toml b/Cargo.toml index 60100a2..e311ea7 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -5,4 +5,5 @@ members = [ "notcuda", "notcuda_inject", "notcuda_redirect", + "ptx", ] \ No newline at end of file diff --git a/notcuda/src/cu.rs b/notcuda/src/cu.rs index fa6f1a5..2b1c827 100644 --- a/notcuda/src/cu.rs +++ b/notcuda/src/cu.rs @@ -166,4 +166,40 @@ impl Context { pub fn null() -> Context { Context(ptr::null_mut()) } +} + +#[repr(transparent)] +pub struct Module(*mut ()); + +#[repr(transparent)] +pub struct Function(*mut ()); + +#[repr(transparent)] +pub struct Stream(*mut ()); + +#[repr(i32)] +#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash, PartialOrd, Ord)] +#[allow(non_camel_case_types)] +pub enum JitOption { + MAX_REGISTERS = 0, + THREADS_PER_BLOCK = 1, + WALL_TIME = 2, + INFO_LOG_BUFFER = 3, + INFO_LOG_BUFFER_SIZE_BYTES = 4, + ERROR_LOG_BUFFER = 5, + ERROR_LOG_BUFFER_SIZE_BYTES = 6, + OPTIMIZATION_LEVEL = 7, + TARGET_FROM_CUCONTEXT = 8, + TARGET = 9, + FALLBACK_STRATEGY = 10, + GENERATE_DEBUG_INFO = 11, + LOG_VERBOSE = 12, + GENERATE_LINE_INFO = 13, + CACHE_MODE = 14, + NEW_SM3X_OPT = 15, + FAST_COMPILE = 16, + GLOBAL_SYMBOL_NAMES = 17, + GLOBAL_SYMBOL_ADDRESSES = 18, + GLOBAL_SYMBOL_COUNT = 19, + NUM_OPTIONS = 20, } \ No newline at end of file diff --git a/notcuda/src/lib.rs b/notcuda/src/lib.rs index fad9f48..707e89e 100644 --- a/notcuda/src/lib.rs +++ b/notcuda/src/lib.rs @@ -223,4 +223,87 @@ pub extern "C" fn cuMemAlloc_v2(dptr: *mut cu::DevicePtr, bytesize: usize) -> cu // TODO: check current context for the device unsafe { l0::zeDriverAllocDeviceMem(drv.base, &descr, bytesize, 0, drv.devices[0].0, dptr as *mut _) } }) +} + +#[no_mangle] +pub extern "C" fn cuCtxDestroy_v2(ctx: cu::Context) -> cu::Result +{ + cu::Result::ERROR_NOT_SUPPORTED +} + +#[no_mangle] +pub extern "C" fn cuMemcpyDtoH_v2(dstHost: *mut (), srcDevice: cu::DevicePtr, byte_count: usize) -> cu::Result +{ + cu::Result::ERROR_NOT_SUPPORTED +} + +#[no_mangle] +pub extern "C" fn cuMemFree_v2(srcDevice: cu::DevicePtr) -> cu::Result +{ + cu::Result::ERROR_NOT_SUPPORTED +} + +#[no_mangle] +pub extern "C" fn cuModuleLoad(module: *mut cu::Module, fname: *const c_char) -> cu::Result +{ + cu::Result::ERROR_NOT_SUPPORTED +} + +#[no_mangle] +pub extern "C" fn cuGetErrorString(error: cu::Result, pStr: *mut *const c_char) -> cu::Result +{ + cu::Result::ERROR_NOT_SUPPORTED +} + +#[no_mangle] +pub extern "C" fn cuLaunchKernel( + f: cu::Function, + gridDimX: c_uint, + gridDimY: c_uint, + gridDimZ: c_uint, + blockDimX: c_uint, + blockDimY: c_uint, + blockDimZ: c_uint, + sharedMemBytes: c_uint, + hStream: cu::Stream, + kernelParams: *mut *mut (), + extra: *mut *mut ()) + -> cu::Result +{ + cu::Result::ERROR_NOT_SUPPORTED +} + +#[no_mangle] +pub extern "C" fn cuModuleLoadDataEx( + module: *mut cu::Module, + image: *const (), + numOptions: c_uint, + options: *mut cu::JitOption, + optionValues: *mut *mut ()) -> cu::Result +{ + cu::Result::ERROR_NOT_SUPPORTED +} + +#[no_mangle] +pub extern "C" fn cuMemcpyHtoD_v2( + dstDevice: cu::DevicePtr, + srcHost: *const (), + ByteCount: usize) -> cu::Result +{ + cu::Result::ERROR_NOT_SUPPORTED +} + +#[no_mangle] +pub extern "C" fn cuCtxCreate_v2(pctx: *mut cu::Context, flags: c_uint, dev: cu::Device) -> cu::Result +{ + cu::Result::SUCCESS +} + +#[no_mangle] +pub extern "C" fn cuModuleGetFunction( + hfunc: *mut cu::Function, + hmod: cu::Module, + name: *const c_char) -> cu::Result +{ + cu::Result::ERROR_NOT_SUPPORTED } \ No newline at end of file diff --git a/ptx/Cargo.toml b/ptx/Cargo.toml new file mode 100644 index 0000000..62050b9 --- /dev/null +++ b/ptx/Cargo.toml @@ -0,0 +1,15 @@ +[package] +name = "ptx" +version = "0.0.0" +authors = ["Andrzej Janik "] +edition = "2018" + +[lib] + +[dependencies] +lalrpop-util = "0.18.1" +regex = "1" + +[build-dependencies.lalrpop] +version = "0.18.1" +features = ["lexer"] diff --git a/ptx/build.rs b/ptx/build.rs new file mode 100644 index 0000000..42c5d59 --- /dev/null +++ b/ptx/build.rs @@ -0,0 +1,5 @@ +extern crate lalrpop; + +fn main() { + lalrpop::process_root().unwrap(); +} \ No newline at end of file diff --git a/ptx/doc/NOTES.md b/ptx/doc/NOTES.md new file mode 100644 index 0000000..8954b4b --- /dev/null +++ b/ptx/doc/NOTES.md @@ -0,0 +1,6 @@ +I'm convinced nobody actually uses parser generators in Rust: +* pomelo can't generate lexer (understandable, as it is a port of lemon and lemon can't do this either) +* pest can't do parse actions, you have to convert your parse tree to ast manually +* lalrpop can't do comments + * and the day I wrote the line above it can +* antlr4rust is untried and requires java to build \ No newline at end of file diff --git a/ptx/src/ast.rs b/ptx/src/ast.rs new file mode 100644 index 0000000..56b299e --- /dev/null +++ b/ptx/src/ast.rs @@ -0,0 +1,18 @@ + +pub struct Module { + version: (u8, u8), + target: Target +} + +pub struct Target { + arch: String, + texturing: TexturingMode, + debug: bool, + f64_to_f32: bool +} + +pub enum TexturingMode { + Unspecified, + Unified, + Independent +} \ No newline at end of file diff --git a/ptx/src/lib.rs b/ptx/src/lib.rs new file mode 100644 index 0000000..7681fcf --- /dev/null +++ b/ptx/src/lib.rs @@ -0,0 +1,16 @@ +#[macro_use] +extern crate lalrpop_util; + +lalrpop_mod!(pub ptx); + +pub mod ast; +pub use ast::Module as Module; + +#[test] +fn version() { + assert!(ptx::ModuleParser::new().parse(" + .version 6.5 + .target + .address_size 64 + ").unwrap() == ()); +} \ No newline at end of file diff --git a/ptx/src/ptx.lalrpop b/ptx/src/ptx.lalrpop new file mode 100644 index 0000000..edc613b --- /dev/null +++ b/ptx/src/ptx.lalrpop @@ -0,0 +1,121 @@ +use std::str::FromStr; +use super::ast; + +grammar; + +match { + r"\s*" => { }, + r"//[^\n\r]*[\n\r]*" => { }, + r"/\*([^\*]*\*+[^\*/])*([^\*]*\*+|[^\*])*\*/" => { }, + _ +} + +pub Module: () = { + Version + Target + Directive* +}; + +Version = { + ".version" U8 "." U8 +}; + +Target = { + ".target" +}; + +Directive : () = { + AddressSize, + Function +}; + +AddressSize = { + ".address_size" "32", + ".address_size" "64" +}; + +Function: (bool, &'input str) = { + LinkingDirective* FunctionInput FunctionBody => (is, id) +}; + +LinkingDirective = { + ".extern", + ".visible", + ".weak" +}; + +IsKernel: bool = { + ".entry" => true, + ".func" => false +}; + +FunctionInput = { + "(" (".param" Type ID)* ")" +}; + +FunctionBody = { + "{" Statement* "}" +}; + +StateSpaceSpecifier = { + ".reg", + ".sreg", + ".const", + ".global", + ".local", + ".param", + ".shared" +}; + +Type = FundamentalType; + +FundamentalType = { + ".s8", ".s16", ".s32", ".s64", + ".u8", ".u16", ".u32", ".u64", + ".f16", ".f16x2", ".f32", ".f64", + ".b8", ".b16", ".b32", ".b64", + ".pred" +}; + +Statement: () = { + Label, + StateSpace, + Instruction +}; + +Label = { + ID ":" +}; + +StateSpace = { + StateSpaceSpecifier Type +}; + +Instruction = { + OpCode Operands ";" +}; + +OpCode = { + "ld.param.u64" +}; + +Operands = Comma; + +Comma: Vec = { // (1) + ",")*> => match e { // (2) + None => v, + Some(e) => { + let mut v = v; + v.push(e); + v + } + } +}; + +Operand = { + ID +}; + +U8: u8 = => u8::from_str(s).unwrap(); + +ID: &'input str = => s; \ No newline at end of file