mirror of
https://github.com/vosen/ZLUDA.git
synced 2025-04-20 00:14:45 +00:00
Parse more source stuff
This commit is contained in:
Andrzej Janik
parent
66e0323c66
commit
0e7338885b
4 changed files with 699 additions and 38 deletions
|
|
@ -3,14 +3,7 @@ extern crate lalrpop_util;
|
|||
|
||||
lalrpop_mod!(pub ptx);
|
||||
|
||||
mod test;
|
||||
|
||||
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() == ());
|
||||
}
|
||||
|
|
@ -4,9 +4,19 @@ use super::ast;
|
|||
grammar;
|
||||
|
||||
match {
|
||||
r"\s*" => { },
|
||||
r"\s+" => { },
|
||||
r"//[^\n\r]*[\n\r]*" => { },
|
||||
r"/\*([^\*]*\*+[^\*/])*([^\*]*\*+|[^\*])*\*/" => { },
|
||||
"ld",
|
||||
"texmode_unified",
|
||||
"texmode_independent",
|
||||
"debug",
|
||||
"map_f64_to_f32",
|
||||
r"sm_[0-9]+" => ShaderModel,
|
||||
} else {
|
||||
r"(?:[a-zA-Z][a-zA-Z0-9_$]*|[_$%][a-zA-Z0-9_$]+)<[0-9]+>" => ParametrizedID,
|
||||
}
|
||||
else {
|
||||
_
|
||||
}
|
||||
|
||||
|
|
@ -17,11 +27,20 @@ pub Module: () = {
|
|||
};
|
||||
|
||||
Version = {
|
||||
".version" U8 "." U8
|
||||
".version" VersionNumber
|
||||
};
|
||||
|
||||
// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#ptx-module-directives-target
|
||||
Target = {
|
||||
".target"
|
||||
".target" Comma<TargetSpecifier>
|
||||
};
|
||||
|
||||
TargetSpecifier = {
|
||||
ShaderModel,
|
||||
"texmode_unified",
|
||||
"texmode_independent",
|
||||
"debug",
|
||||
"map_f64_to_f32"
|
||||
};
|
||||
|
||||
Directive : () = {
|
||||
|
|
@ -35,7 +54,7 @@ AddressSize = {
|
|||
};
|
||||
|
||||
Function: (bool, &'input str) = {
|
||||
LinkingDirective* <is:IsKernel> <id:ID> FunctionInput FunctionBody => (is, id)
|
||||
LinkingDirective* <is:IsKernel> <id:ID> "(" Comma<FunctionInput> ")" FunctionBody => (is, id)
|
||||
};
|
||||
|
||||
LinkingDirective = {
|
||||
|
|
@ -50,7 +69,7 @@ IsKernel: bool = {
|
|||
};
|
||||
|
||||
FunctionInput = {
|
||||
"(" (".param" Type ID)* ")"
|
||||
".param" Type ID
|
||||
};
|
||||
|
||||
FunctionBody = {
|
||||
|
|
@ -63,46 +82,91 @@ StateSpaceSpecifier = {
|
|||
".const",
|
||||
".global",
|
||||
".local",
|
||||
".param",
|
||||
".shared"
|
||||
};
|
||||
|
||||
Type = FundamentalType;
|
||||
Type = {
|
||||
BaseType,
|
||||
".pred",
|
||||
".f16",
|
||||
".f16x2",
|
||||
};
|
||||
|
||||
FundamentalType = {
|
||||
".s8", ".s16", ".s32", ".s64",
|
||||
".u8", ".u16", ".u32", ".u64",
|
||||
".f16", ".f16x2", ".f32", ".f64",
|
||||
BaseType = {
|
||||
".b8", ".b16", ".b32", ".b64",
|
||||
".pred"
|
||||
".u8", ".u16", ".u32", ".u64",
|
||||
".s8", ".s16", ".s32", ".s64",
|
||||
".f32", ".f64"
|
||||
};
|
||||
|
||||
Statement: () = {
|
||||
Label,
|
||||
StateSpace,
|
||||
Instruction
|
||||
DebugDirective,
|
||||
Variable ";",
|
||||
Instruction ";"
|
||||
};
|
||||
|
||||
DebugDirective = {
|
||||
DebugLocation
|
||||
};
|
||||
|
||||
// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#debugging-directives-loc
|
||||
DebugLocation = {
|
||||
".loc" Num Num Num
|
||||
};
|
||||
|
||||
Label = {
|
||||
ID ":"
|
||||
};
|
||||
|
||||
StateSpace = {
|
||||
StateSpaceSpecifier Type
|
||||
Variable = {
|
||||
StateSpaceSpecifier Type VariableName
|
||||
};
|
||||
|
||||
VariableName = {
|
||||
ID,
|
||||
ParametrizedID
|
||||
};
|
||||
|
||||
Instruction = {
|
||||
OpCode Operands ";"
|
||||
InstLd
|
||||
};
|
||||
|
||||
OpCode = {
|
||||
"ld.param.u64"
|
||||
// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-ld
|
||||
InstLd = {
|
||||
"ld" LdQualifier? LdStateSpace? LdCacheOperator? Vector? BaseType ID "," "[" ID "]"
|
||||
};
|
||||
|
||||
Operands = Comma<Operand>;
|
||||
LdQualifier = {
|
||||
".weak",
|
||||
".volatile",
|
||||
".relaxed.scope",
|
||||
".acquire.scope",
|
||||
};
|
||||
|
||||
Comma<T>: Vec<T> = { // (1)
|
||||
<v:(<T> ",")*> <e:T?> => match e { // (2)
|
||||
LdStateSpace = {
|
||||
".const",
|
||||
".global",
|
||||
".local",
|
||||
".param",
|
||||
".shared",
|
||||
};
|
||||
|
||||
LdCacheOperator = {
|
||||
".ca",
|
||||
".cg",
|
||||
".cs",
|
||||
".lu",
|
||||
".cv",
|
||||
};
|
||||
|
||||
Vector = {
|
||||
".v2",
|
||||
".v4"
|
||||
};
|
||||
|
||||
Comma<T>: Vec<T> = {
|
||||
<v:(<T> ",")*> <e:T?> => match e {
|
||||
None => v,
|
||||
Some(e) => {
|
||||
let mut v = v;
|
||||
|
|
@ -112,10 +176,7 @@ Comma<T>: Vec<T> = { // (1)
|
|||
}
|
||||
};
|
||||
|
||||
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;
|
||||
VersionNumber = r"[0-9]+\.[0-9]+";
|
||||
Num: u64 = <s:r"[0-9]+"> => u64::from_str(s).unwrap();
|
||||
// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#identifiers
|
||||
ID: &'input str = <s:r"[a-zA-Z][a-zA-Z0-9_$]*|[_$%][a-zA-Z0-9_$]+"> => s;
|
||||
15
ptx/src/test/mod.rs
Normal file
15
ptx/src/test/mod.rs
Normal file
|
|
@ -0,0 +1,15 @@
|
|||
use super::ptx;
|
||||
|
||||
|
||||
#[test]
|
||||
fn empty() {
|
||||
assert!(ptx::ModuleParser::new().parse(
|
||||
".version 6.5 .target sm_30, debug")
|
||||
.unwrap() == ());
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn vector_add() {
|
||||
let vector_add = include_str!("vectorAdd_kernel64.ptx");
|
||||
assert!(ptx::ModuleParser::new().parse(vector_add).unwrap() == ());
|
||||
}
|
||||
592
ptx/src/test/vectorAdd_kernel64.ptx
Normal file
592
ptx/src/test/vectorAdd_kernel64.ptx
Normal file
|
|
@ -0,0 +1,592 @@
|
|||
//
|
||||
// Generated by NVIDIA NVVM Compiler
|
||||
//
|
||||
// Compiler Build ID: CL-27506705
|
||||
// Cuda compilation tools, release 10.2, V10.2.89
|
||||
// Based on LLVM 3.4svn
|
||||
//
|
||||
|
||||
.version 6.5
|
||||
.target sm_30, debug
|
||||
.address_size 64
|
||||
|
||||
// .globl VecAdd_kernel
|
||||
|
||||
.visible .entry VecAdd_kernel(
|
||||
.param .u64 VecAdd_kernel_param_0,
|
||||
.param .u64 VecAdd_kernel_param_1,
|
||||
.param .u64 VecAdd_kernel_param_2,
|
||||
.param .u32 VecAdd_kernel_param_3
|
||||
)
|
||||
{
|
||||
.reg .pred %p<3>;
|
||||
.reg .f32 %f<4>;
|
||||
.reg .b32 %r<7>;
|
||||
.reg .b64 %rd<13>;
|
||||
|
||||
|
||||
.loc 1 21 1
|
||||
func_begin0:
|
||||
.loc 1 0 0
|
||||
|
||||
.loc 1 21 1
|
||||
|
||||
ld.param.u64 %rd1, [VecAdd_kernel_param_0];
|
||||
ld.param.u64 %rd2, [VecAdd_kernel_param_1];
|
||||
ld.param.u64 %rd3, [VecAdd_kernel_param_2];
|
||||
ld.param.u32 %r2, [VecAdd_kernel_param_3];
|
||||
func_exec_begin0:
|
||||
.loc 1 23 11
|
||||
tmp0:
|
||||
mov.u32 %r3, %ntid.x;
|
||||
mov.u32 %r4, %ctaid.x;
|
||||
mul.lo.s32 %r5, %r3, %r4;
|
||||
mov.u32 %r6, %tid.x;
|
||||
add.s32 %r1, %r5, %r6;
|
||||
tmp1:
|
||||
.loc 1 25 5
|
||||
setp.lt.s32 %p1, %r1, %r2;
|
||||
not.pred %p2, %p1;
|
||||
@%p2 bra BB0_2;
|
||||
bra.uni BB0_1;
|
||||
|
||||
BB0_1:
|
||||
.loc 1 26 9
|
||||
tmp2:
|
||||
cvt.s64.s32 %rd4, %r1;
|
||||
shl.b64 %rd5, %rd4, 2;
|
||||
add.s64 %rd6, %rd1, %rd5;
|
||||
ld.f32 %f1, [%rd6];
|
||||
cvt.s64.s32 %rd7, %r1;
|
||||
shl.b64 %rd8, %rd7, 2;
|
||||
add.s64 %rd9, %rd2, %rd8;
|
||||
ld.f32 %f2, [%rd9];
|
||||
add.f32 %f3, %f1, %f2;
|
||||
cvt.s64.s32 %rd10, %r1;
|
||||
shl.b64 %rd11, %rd10, 2;
|
||||
add.s64 %rd12, %rd3, %rd11;
|
||||
st.f32 [%rd12], %f3;
|
||||
tmp3:
|
||||
|
||||
BB0_2:
|
||||
.loc 1 27 1
|
||||
ret;
|
||||
tmp4:
|
||||
func_end0:
|
||||
}
|
||||
|
||||
.file 1 "/home/vosen/cuda-samples/0_Simple/vectorAddMMAP/vectorAdd_kernel.cu", 1581801938, 860
|
||||
|
||||
.section .debug_info {
|
||||
.b32 314
|
||||
.b8 2
|
||||
.b8 0
|
||||
.b32 .debug_abbrev
|
||||
.b8 8
|
||||
.b8 1
|
||||
|
||||
.b8 108
|
||||
.b8 103
|
||||
.b8 101
|
||||
.b8 110
|
||||
.b8 102
|
||||
.b8 101
|
||||
.b8 58
|
||||
.b8 32
|
||||
.b8 69
|
||||
.b8 68
|
||||
.b8 71
|
||||
.b8 32
|
||||
.b8 53
|
||||
.b8 46
|
||||
.b8 48
|
||||
|
||||
.b8 0
|
||||
.b8 4
|
||||
.b8 118
|
||||
.b8 101
|
||||
.b8 99
|
||||
.b8 116
|
||||
.b8 111
|
||||
.b8 114
|
||||
.b8 65
|
||||
.b8 100
|
||||
.b8 100
|
||||
.b8 95
|
||||
.b8 107
|
||||
.b8 101
|
||||
.b8 114
|
||||
.b8 110
|
||||
.b8 101
|
||||
.b8 108
|
||||
.b8 46
|
||||
.b8 99
|
||||
.b8 117
|
||||
|
||||
.b8 0
|
||||
.b64 0
|
||||
.b32 .debug_line
|
||||
.b8 47
|
||||
.b8 104
|
||||
.b8 111
|
||||
.b8 109
|
||||
.b8 101
|
||||
.b8 47
|
||||
.b8 118
|
||||
.b8 111
|
||||
.b8 115
|
||||
.b8 101
|
||||
.b8 110
|
||||
.b8 47
|
||||
.b8 99
|
||||
.b8 117
|
||||
.b8 100
|
||||
.b8 97
|
||||
.b8 45
|
||||
.b8 115
|
||||
.b8 97
|
||||
.b8 109
|
||||
.b8 112
|
||||
.b8 108
|
||||
.b8 101
|
||||
.b8 115
|
||||
.b8 47
|
||||
.b8 48
|
||||
.b8 95
|
||||
.b8 83
|
||||
.b8 105
|
||||
.b8 109
|
||||
.b8 112
|
||||
.b8 108
|
||||
.b8 101
|
||||
.b8 47
|
||||
.b8 118
|
||||
.b8 101
|
||||
.b8 99
|
||||
.b8 116
|
||||
.b8 111
|
||||
.b8 114
|
||||
.b8 65
|
||||
.b8 100
|
||||
.b8 100
|
||||
.b8 77
|
||||
.b8 77
|
||||
.b8 65
|
||||
.b8 80
|
||||
|
||||
.b8 0
|
||||
.b8 2
|
||||
|
||||
.b8 86
|
||||
.b8 101
|
||||
.b8 99
|
||||
.b8 65
|
||||
.b8 100
|
||||
.b8 100
|
||||
.b8 95
|
||||
.b8 107
|
||||
.b8 101
|
||||
.b8 114
|
||||
.b8 110
|
||||
.b8 101
|
||||
.b8 108
|
||||
|
||||
.b8 0
|
||||
.b8 86
|
||||
.b8 101
|
||||
.b8 99
|
||||
.b8 65
|
||||
.b8 100
|
||||
.b8 100
|
||||
.b8 95
|
||||
.b8 107
|
||||
.b8 101
|
||||
.b8 114
|
||||
.b8 110
|
||||
.b8 101
|
||||
.b8 108
|
||||
|
||||
.b8 0
|
||||
.b8 1
|
||||
.b8 21
|
||||
.b32 278
|
||||
.b8 1
|
||||
.b64 func_begin0
|
||||
.b64 func_end0
|
||||
.b8 1
|
||||
.b8 156
|
||||
.b8 3
|
||||
|
||||
.b8 65
|
||||
|
||||
.b8 0
|
||||
.b8 1
|
||||
.b8 21
|
||||
.b32 284
|
||||
.b8 9
|
||||
.b8 3
|
||||
.b64 VecAdd_kernel_param_0
|
||||
.b8 7
|
||||
.b8 3
|
||||
|
||||
.b8 66
|
||||
|
||||
.b8 0
|
||||
.b8 1
|
||||
.b8 21
|
||||
.b32 284
|
||||
.b8 9
|
||||
.b8 3
|
||||
.b64 VecAdd_kernel_param_1
|
||||
.b8 7
|
||||
.b8 3
|
||||
|
||||
.b8 67
|
||||
|
||||
.b8 0
|
||||
.b8 1
|
||||
.b8 21
|
||||
.b32 304
|
||||
.b8 9
|
||||
.b8 3
|
||||
.b64 VecAdd_kernel_param_2
|
||||
.b8 7
|
||||
.b8 3
|
||||
|
||||
.b8 78
|
||||
|
||||
.b8 0
|
||||
.b8 1
|
||||
.b8 21
|
||||
.b32 310
|
||||
.b8 9
|
||||
.b8 3
|
||||
.b64 VecAdd_kernel_param_3
|
||||
.b8 7
|
||||
.b8 4
|
||||
|
||||
.b64 tmp0
|
||||
.b64 tmp4
|
||||
.b8 5
|
||||
|
||||
.b8 105
|
||||
|
||||
.b8 0
|
||||
.b8 1
|
||||
.b8 23
|
||||
.b32 310
|
||||
.b8 5
|
||||
.b8 144
|
||||
.b8 177
|
||||
.b8 228
|
||||
.b8 149
|
||||
.b8 1
|
||||
.b8 2
|
||||
.b8 0
|
||||
.b8 0
|
||||
.b8 6
|
||||
|
||||
.b8 118
|
||||
.b8 111
|
||||
.b8 105
|
||||
.b8 100
|
||||
|
||||
.b8 0
|
||||
.b8 7
|
||||
|
||||
.b32 290
|
||||
.b8 12
|
||||
.b8 8
|
||||
|
||||
.b32 295
|
||||
.b8 9
|
||||
|
||||
.b8 102
|
||||
.b8 108
|
||||
.b8 111
|
||||
.b8 97
|
||||
.b8 116
|
||||
|
||||
.b8 0
|
||||
.b8 4
|
||||
.b8 4
|
||||
.b8 7
|
||||
|
||||
.b32 295
|
||||
.b8 12
|
||||
.b8 9
|
||||
|
||||
.b8 105
|
||||
.b8 110
|
||||
.b8 116
|
||||
|
||||
.b8 0
|
||||
.b8 5
|
||||
.b8 4
|
||||
.b8 0
|
||||
}
|
||||
.section .debug_abbrev {
|
||||
.b8 1
|
||||
|
||||
.b8 17
|
||||
|
||||
.b8 1
|
||||
|
||||
.b8 37
|
||||
|
||||
.b8 8
|
||||
|
||||
.b8 19
|
||||
|
||||
.b8 11
|
||||
|
||||
.b8 3
|
||||
|
||||
.b8 8
|
||||
|
||||
.b8 17
|
||||
|
||||
.b8 1
|
||||
|
||||
.b8 16
|
||||
|
||||
.b8 6
|
||||
|
||||
.b8 27
|
||||
|
||||
.b8 8
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 2
|
||||
|
||||
.b8 46
|
||||
|
||||
.b8 1
|
||||
|
||||
.b8 135
|
||||
.b8 64
|
||||
|
||||
.b8 8
|
||||
|
||||
.b8 3
|
||||
|
||||
.b8 8
|
||||
|
||||
.b8 58
|
||||
|
||||
.b8 11
|
||||
|
||||
.b8 59
|
||||
|
||||
.b8 11
|
||||
|
||||
.b8 73
|
||||
|
||||
.b8 19
|
||||
|
||||
.b8 63
|
||||
|
||||
.b8 12
|
||||
|
||||
.b8 17
|
||||
|
||||
.b8 1
|
||||
|
||||
.b8 18
|
||||
|
||||
.b8 1
|
||||
|
||||
.b8 64
|
||||
|
||||
.b8 10
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 3
|
||||
|
||||
.b8 5
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 3
|
||||
|
||||
.b8 8
|
||||
|
||||
.b8 58
|
||||
|
||||
.b8 11
|
||||
|
||||
.b8 59
|
||||
|
||||
.b8 11
|
||||
|
||||
.b8 73
|
||||
|
||||
.b8 19
|
||||
|
||||
.b8 2
|
||||
|
||||
.b8 10
|
||||
|
||||
.b8 51
|
||||
|
||||
.b8 11
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 4
|
||||
|
||||
.b8 11
|
||||
|
||||
.b8 1
|
||||
|
||||
.b8 17
|
||||
|
||||
.b8 1
|
||||
|
||||
.b8 18
|
||||
|
||||
.b8 1
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 5
|
||||
|
||||
.b8 52
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 3
|
||||
|
||||
.b8 8
|
||||
|
||||
.b8 58
|
||||
|
||||
.b8 11
|
||||
|
||||
.b8 59
|
||||
|
||||
.b8 11
|
||||
|
||||
.b8 73
|
||||
|
||||
.b8 19
|
||||
|
||||
.b8 2
|
||||
|
||||
.b8 10
|
||||
|
||||
.b8 51
|
||||
|
||||
.b8 11
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 6
|
||||
|
||||
.b8 59
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 3
|
||||
|
||||
.b8 8
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 7
|
||||
|
||||
.b8 15
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 73
|
||||
|
||||
.b8 19
|
||||
|
||||
.b8 51
|
||||
|
||||
.b8 11
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 8
|
||||
|
||||
.b8 38
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 73
|
||||
|
||||
.b8 19
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 9
|
||||
|
||||
.b8 36
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 3
|
||||
|
||||
.b8 8
|
||||
|
||||
.b8 62
|
||||
|
||||
.b8 11
|
||||
|
||||
.b8 11
|
||||
|
||||
.b8 11
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 0
|
||||
|
||||
.b8 0
|
||||
|
||||
}
|
||||
.section .debug_ranges {
|
||||
}
|
||||
.section .debug_pubnames {
|
||||
.b32 32
|
||||
.b8 2
|
||||
.b8 0
|
||||
.b32 .debug_info
|
||||
.b32 314
|
||||
.b32 109
|
||||
.b8 86
|
||||
.b8 101
|
||||
.b8 99
|
||||
.b8 65
|
||||
.b8 100
|
||||
.b8 100
|
||||
.b8 95
|
||||
.b8 107
|
||||
.b8 101
|
||||
.b8 114
|
||||
.b8 110
|
||||
.b8 101
|
||||
.b8 108
|
||||
.b8 0
|
||||
|
||||
.b32 0
|
||||
}
|
||||
Loading…
Add table
Reference in a new issue