Start parsing PTX text

Cargo.toml

@@ -5,4 +5,5 @@ members = [
     "notcuda",
     "notcuda_inject",
     "notcuda_redirect",
+    "ptx",
 ]

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,
 }

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
 }

ptx/Cargo.toml

@@ -0,0 +1,15 @@
+[package]
+name = "ptx"
+version = "0.0.0"
+authors = ["Andrzej Janik <vosen@vosen.pl>"]
+edition = "2018"
+
+[lib]
+
+[dependencies]
+lalrpop-util = "0.18.1"
+regex = "1"
+
+[build-dependencies.lalrpop]
+version = "0.18.1"
+features = ["lexer"]

ptx/build.rs

@@ -0,0 +1,5 @@
+extern crate lalrpop;
+
+fn main() {
+    lalrpop::process_root().unwrap();
+}

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

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
+}

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() == ());
+}

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* <is:IsKernel> <id:ID> 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<Operand>;
+
+Comma<T>: Vec<T> = { // (1)
+    <v:(<T> ",")*> <e:T?> => match e { // (2)
+        None => v,
+        Some(e) => {
+            let mut v = v;
+            v.push(e);
+            v
+        }
+    }
+};
+
+Operand = {
+    ID
+};
+
+U8: u8 = <s:r"[0-9]+"> => u8::from_str(s).unwrap();
+
+ID: &'input str = <s:"[a-zA-Z][a-zA-Z0-9_$]*|[_$%][a-zA-Z0-9_$]"> => s;