aboutsummaryrefslogtreecommitdiffhomepage
path: root/ptx
diff options
context:
space:
mode:
Diffstat (limited to 'ptx')
-rw-r--r--ptx/src/test/spirv_run/add_tuning.spvtxt103
-rw-r--r--ptx/src/test/spirv_run/extern_func.ptx39
-rw-r--r--ptx/src/test/spirv_run/extern_func.spvtxt75
-rw-r--r--ptx/src/test/spirv_run/func_ptr.spvtxt60
-rw-r--r--ptx/src/test/spirv_run/mod.rs17
-rw-r--r--ptx/src/translate.rs88
6 files changed, 270 insertions, 112 deletions
diff --git a/ptx/src/test/spirv_run/add_tuning.spvtxt b/ptx/src/test/spirv_run/add_tuning.spvtxt
index 173e0d4..d65f04d 100644
--- a/ptx/src/test/spirv_run/add_tuning.spvtxt
+++ b/ptx/src/test/spirv_run/add_tuning.spvtxt
@@ -1,48 +1,55 @@
- OpCapability GenericPointer
- OpCapability Linkage
- OpCapability Addresses
- OpCapability Kernel
- OpCapability Int8
- OpCapability Int16
- OpCapability Int64
- OpCapability Float16
- OpCapability Float64
- %23 = OpExtInstImport "OpenCL.std"
- OpMemoryModel Physical64 OpenCL
- OpEntryPoint Kernel %1 "add_tuning"
- OpExecutionMode %1 MaxWorkgroupSizeINTEL 256 1 1
- %void = OpTypeVoid
- %ulong = OpTypeInt 64 0
- %26 = OpTypeFunction %void %ulong %ulong
-%_ptr_Function_ulong = OpTypePointer Function %ulong
-%_ptr_Generic_ulong = OpTypePointer Generic %ulong
- %ulong_1 = OpConstant %ulong 1
- %1 = OpFunction %void None %26
- %8 = OpFunctionParameter %ulong
- %9 = OpFunctionParameter %ulong
- %21 = OpLabel
- %2 = OpVariable %_ptr_Function_ulong Function
- %3 = OpVariable %_ptr_Function_ulong Function
- %4 = OpVariable %_ptr_Function_ulong Function
- %5 = OpVariable %_ptr_Function_ulong Function
- %6 = OpVariable %_ptr_Function_ulong Function
- %7 = OpVariable %_ptr_Function_ulong Function
- OpStore %2 %8
- OpStore %3 %9
- %10 = OpLoad %ulong %2 Aligned 8
- OpStore %4 %10
- %11 = OpLoad %ulong %3 Aligned 8
- OpStore %5 %11
- %13 = OpLoad %ulong %4
- %19 = OpConvertUToPtr %_ptr_Generic_ulong %13
- %12 = OpLoad %ulong %19 Aligned 8
- OpStore %6 %12
- %15 = OpLoad %ulong %6
- %14 = OpIAdd %ulong %15 %ulong_1
- OpStore %7 %14
- %16 = OpLoad %ulong %5
- %17 = OpLoad %ulong %7
- %20 = OpConvertUToPtr %_ptr_Generic_ulong %16
- OpStore %20 %17 Aligned 8
- OpReturn
- OpFunctionEnd
+; SPIR-V
+; Version: 1.3
+; Generator: rspirv
+; Bound: 29
+OpCapability GenericPointer
+OpCapability Linkage
+OpCapability Addresses
+OpCapability Kernel
+OpCapability Int8
+OpCapability Int16
+OpCapability Int64
+OpCapability Float16
+OpCapability Float64
+OpCapability DenormFlushToZero
+%23 = OpExtInstImport "OpenCL.std"
+OpMemoryModel Physical64 OpenCL
+OpEntryPoint Kernel %1 "add_tuning"
+OpExecutionMode %1 ContractionOff
+; OpExecutionMode %1 MaxWorkgroupSizeINTEL 256 1 1
+OpDecorate %1 LinkageAttributes "add_tuning" Export
+%24 = OpTypeVoid
+%25 = OpTypeInt 64 0
+%26 = OpTypeFunction %24 %25 %25
+%27 = OpTypePointer Function %25
+%28 = OpTypePointer Generic %25
+%18 = OpConstant %25 1
+%1 = OpFunction %24 None %26
+%8 = OpFunctionParameter %25
+%9 = OpFunctionParameter %25
+%21 = OpLabel
+%2 = OpVariable %27 Function
+%3 = OpVariable %27 Function
+%4 = OpVariable %27 Function
+%5 = OpVariable %27 Function
+%6 = OpVariable %27 Function
+%7 = OpVariable %27 Function
+OpStore %2 %8
+OpStore %3 %9
+%10 = OpLoad %25 %2 Aligned 8
+OpStore %4 %10
+%11 = OpLoad %25 %3 Aligned 8
+OpStore %5 %11
+%13 = OpLoad %25 %4
+%19 = OpConvertUToPtr %28 %13
+%12 = OpLoad %25 %19 Aligned 8
+OpStore %6 %12
+%15 = OpLoad %25 %6
+%14 = OpIAdd %25 %15 %18
+OpStore %7 %14
+%16 = OpLoad %25 %5
+%17 = OpLoad %25 %7
+%20 = OpConvertUToPtr %28 %16
+OpStore %20 %17 Aligned 8
+OpReturn
+OpFunctionEnd \ No newline at end of file
diff --git a/ptx/src/test/spirv_run/extern_func.ptx b/ptx/src/test/spirv_run/extern_func.ptx
new file mode 100644
index 0000000..5e3a4ce
--- /dev/null
+++ b/ptx/src/test/spirv_run/extern_func.ptx
@@ -0,0 +1,39 @@
+.version 6.5
+.target sm_30
+.address_size 64
+
+.extern .func (.param .align 16 .b8 func_retval0[16]) foobar
+(
+ .param .b64 extern_func
+)
+;
+
+.visible .entry extern_func(
+ .param .u64 input,
+ .param .u64 output
+)
+{
+ .reg .u64 in_addr;
+ .reg .u64 out_addr;
+ .reg .u64 temp;
+ .reg .u64 temp2;
+
+ ld.param.u64 in_addr, [input];
+ ld.param.u64 out_addr, [output];
+
+ ld.global.u64 temp, [in_addr];
+
+ {
+ .param .b64 param0;
+ st.param.b64 [param0+0], temp;
+ .param .align 16 .b8 retval0[16];
+ call.uni (retval0) ,
+ foobar ,
+ (
+ param0
+ );
+ ld.param.u64 temp2, [retval0];
+ }
+ st.u64 [out_addr], temp2;
+ ret;
+}
diff --git a/ptx/src/test/spirv_run/extern_func.spvtxt b/ptx/src/test/spirv_run/extern_func.spvtxt
new file mode 100644
index 0000000..b757029
--- /dev/null
+++ b/ptx/src/test/spirv_run/extern_func.spvtxt
@@ -0,0 +1,75 @@
+ OpCapability GenericPointer
+ OpCapability Linkage
+ OpCapability Addresses
+ OpCapability Kernel
+ OpCapability Int8
+ OpCapability Int16
+ OpCapability Int64
+ OpCapability Float16
+ OpCapability Float64
+ OpCapability DenormFlushToZero
+ %31 = OpExtInstImport "OpenCL.std"
+ OpMemoryModel Physical64 OpenCL
+ OpEntryPoint Kernel %4 "extern_func"
+ OpExecutionMode %4 ContractionOff
+ OpDecorate %1 LinkageAttributes "foobar" Import
+ OpDecorate %12 Alignment 16
+ OpDecorate %4 LinkageAttributes "extern_func" Export
+ %void = OpTypeVoid
+ %ulong = OpTypeInt 64 0
+%_ptr_Function_ulong = OpTypePointer Function %ulong
+ %uint = OpTypeInt 32 0
+ %uchar = OpTypeInt 8 0
+ %uint_16 = OpConstant %uint 16
+%_arr_uchar_uint_16 = OpTypeArray %uchar %uint_16
+%_ptr_Function__arr_uchar_uint_16 = OpTypePointer Function %_arr_uchar_uint_16
+ %40 = OpTypeFunction %void %_ptr_Function_ulong %_ptr_Function__arr_uchar_uint_16
+ %uint_16_0 = OpConstant %uint 16
+ %42 = OpTypeFunction %void %ulong %ulong
+ %uint_16_1 = OpConstant %uint 16
+%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
+ %ulong_0 = OpConstant %ulong 0
+%_ptr_Function_uchar = OpTypePointer Function %uchar
+%_ptr_Generic_ulong = OpTypePointer Generic %ulong
+ %1 = OpFunction %void None %40
+ %3 = OpFunctionParameter %_ptr_Function_ulong
+ %2 = OpFunctionParameter %_ptr_Function__arr_uchar_uint_16
+ OpFunctionEnd
+ %4 = OpFunction %void None %42
+ %13 = OpFunctionParameter %ulong
+ %14 = OpFunctionParameter %ulong
+ %29 = OpLabel
+ %5 = OpVariable %_ptr_Function_ulong Function
+ %6 = OpVariable %_ptr_Function_ulong Function
+ %7 = OpVariable %_ptr_Function_ulong Function
+ %8 = OpVariable %_ptr_Function_ulong Function
+ %9 = OpVariable %_ptr_Function_ulong Function
+ %10 = OpVariable %_ptr_Function_ulong Function
+ %11 = OpVariable %_ptr_Function_ulong Function
+ %12 = OpVariable %_ptr_Function__arr_uchar_uint_16 Function
+ OpStore %5 %13
+ OpStore %6 %14
+ %15 = OpLoad %ulong %5 Aligned 8
+ OpStore %7 %15
+ %16 = OpLoad %ulong %6 Aligned 8
+ OpStore %8 %16
+ %18 = OpLoad %ulong %7
+ %25 = OpConvertUToPtr %_ptr_CrossWorkgroup_ulong %18
+ %17 = OpLoad %ulong %25 Aligned 8
+ OpStore %9 %17
+ %19 = OpLoad %ulong %9
+ %46 = OpBitcast %_ptr_Function_uchar %11
+ %47 = OpInBoundsPtrAccessChain %_ptr_Function_uchar %46 %ulong_0
+ %24 = OpBitcast %_ptr_Function_ulong %47
+ %26 = OpCopyObject %ulong %19
+ OpStore %24 %26 Aligned 8
+ %48 = OpFunctionCall %void %1 %11 %12
+ %27 = OpBitcast %_ptr_Function_ulong %12
+ %20 = OpLoad %ulong %27 Aligned 8
+ OpStore %10 %20
+ %21 = OpLoad %ulong %8
+ %22 = OpLoad %ulong %10
+ %28 = OpConvertUToPtr %_ptr_Generic_ulong %21
+ OpStore %28 %22 Aligned 8
+ OpReturn
+ OpFunctionEnd \ No newline at end of file
diff --git a/ptx/src/test/spirv_run/func_ptr.spvtxt b/ptx/src/test/spirv_run/func_ptr.spvtxt
index adc71eb..4ff74c6 100644
--- a/ptx/src/test/spirv_run/func_ptr.spvtxt
+++ b/ptx/src/test/spirv_run/func_ptr.spvtxt
@@ -7,24 +7,26 @@
OpCapability Int64
OpCapability Float16
OpCapability Float64
- %38 = OpExtInstImport "OpenCL.std"
+ OpCapability DenormFlushToZero
+ %39 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
- OpEntryPoint Kernel %11 "func_ptr"
- OpExecutionMode %11 ContractionOff
+ OpEntryPoint Kernel %12 "func_ptr"
+ OpExecutionMode %12 ContractionOff
+ OpDecorate %12 LinkageAttributes "func_ptr" Export
%void = OpTypeVoid
%float = OpTypeFloat 32
- %41 = OpTypeFunction %float %float %float
+ %42 = OpTypeFunction %float %float %float
%_ptr_Function_float = OpTypePointer Function %float
%ulong = OpTypeInt 64 0
- %44 = OpTypeFunction %void %ulong %ulong
+ %45 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Generic_ulong = OpTypePointer Generic %ulong
%ulong_1 = OpConstant %ulong 1
%ulong_0 = OpConstant %ulong 0
- %1 = OpFunction %float None %41
+ %1 = OpFunction %float None %42
%5 = OpFunctionParameter %float
%6 = OpFunctionParameter %float
- %10 = OpLabel
+ %11 = OpLabel
%3 = OpVariable %_ptr_Function_float Function
%4 = OpVariable %_ptr_Function_float Function
%2 = OpVariable %_ptr_Function_float Function
@@ -34,40 +36,42 @@
%9 = OpLoad %float %4
%7 = OpFAdd %float %8 %9
OpStore %2 %7
+ %10 = OpLoad %float %2
+ OpReturnValue %10
OpFunctionEnd
- %11 = OpFunction %void None %44
- %19 = OpFunctionParameter %ulong
+ %12 = OpFunction %void None %45
%20 = OpFunctionParameter %ulong
- %36 = OpLabel
- %12 = OpVariable %_ptr_Function_ulong Function
+ %21 = OpFunctionParameter %ulong
+ %37 = OpLabel
%13 = OpVariable %_ptr_Function_ulong Function
%14 = OpVariable %_ptr_Function_ulong Function
%15 = OpVariable %_ptr_Function_ulong Function
%16 = OpVariable %_ptr_Function_ulong Function
%17 = OpVariable %_ptr_Function_ulong Function
%18 = OpVariable %_ptr_Function_ulong Function
- OpStore %12 %19
+ %19 = OpVariable %_ptr_Function_ulong Function
OpStore %13 %20
- %21 = OpLoad %ulong %12 Aligned 8
OpStore %14 %21
%22 = OpLoad %ulong %13 Aligned 8
OpStore %15 %22
- %24 = OpLoad %ulong %14
- %34 = OpConvertUToPtr %_ptr_Generic_ulong %24
- %23 = OpLoad %ulong %34 Aligned 8
+ %23 = OpLoad %ulong %14 Aligned 8
OpStore %16 %23
- %26 = OpLoad %ulong %16
- %25 = OpIAdd %ulong %26 %ulong_1
- OpStore %17 %25
- %27 = OpCopyObject %ulong %ulong_0
- OpStore %18 %27
- %29 = OpLoad %ulong %17
+ %25 = OpLoad %ulong %15
+ %35 = OpConvertUToPtr %_ptr_Generic_ulong %25
+ %24 = OpLoad %ulong %35 Aligned 8
+ OpStore %17 %24
+ %27 = OpLoad %ulong %17
+ %26 = OpIAdd %ulong %27 %ulong_1
+ OpStore %18 %26
+ %28 = OpCopyObject %ulong %ulong_0
+ OpStore %19 %28
%30 = OpLoad %ulong %18
- %28 = OpIAdd %ulong %29 %30
- OpStore %17 %28
- %31 = OpLoad %ulong %15
- %32 = OpLoad %ulong %17
- %35 = OpConvertUToPtr %_ptr_Generic_ulong %31
- OpStore %35 %32 Aligned 8
+ %31 = OpLoad %ulong %19
+ %29 = OpIAdd %ulong %30 %31
+ OpStore %18 %29
+ %32 = OpLoad %ulong %16
+ %33 = OpLoad %ulong %18
+ %36 = OpConvertUToPtr %_ptr_Generic_ulong %32
+ OpStore %36 %33 Aligned 8
OpReturn
OpFunctionEnd
diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs
index d68cf17..be34d0f 100644
--- a/ptx/src/test/spirv_run/mod.rs
+++ b/ptx/src/test/spirv_run/mod.rs
@@ -60,6 +60,18 @@ macro_rules! test_ptx {
}
}
};
+
+ ($fn_name:ident) => {
+ paste::item! {
+ #[test]
+ fn [<$fn_name _spvtxt>]() -> Result<(), Box<dyn std::error::Error>> {
+ let ptx_txt = include_str!(concat!(stringify!($fn_name), ".ptx"));
+ let spirv_file_name = concat!(stringify!($fn_name), ".spvtxt");
+ let spirv_txt = include_bytes!(concat!(stringify!($fn_name), ".spvtxt"));
+ test_spvtxt_assert(ptx_txt, spirv_txt, spirv_file_name)
+ }
+ }
+ };
}
test_ptx!(ld_st, [1u64], [1u64]);
@@ -209,8 +221,9 @@ test_ptx!(cvt_f64_f32, [0.125f32], [0.125f64]);
test_ptx!(prmt, [0x70c507d6u32, 0x6fbd4b5cu32], [0x6fbdd65cu32]);
test_ptx!(activemask, [0u32], [1u32]);
test_ptx!(membar, [152731u32], [152731u32]);
-test_ptx!(func_ptr, [152731u64], [152732u64]);
-test_ptx!(lanemask_lt, [187235u32], [187236u32]);
+test_ptx!(func_ptr);
+test_ptx!(lanemask_lt);
+test_ptx!(extern_func);
struct DisplayError<T: Debug> {
err: T,
diff --git a/ptx/src/translate.rs b/ptx/src/translate.rs
index 297588a..13c578b 100644
--- a/ptx/src/translate.rs
+++ b/ptx/src/translate.rs
@@ -543,10 +543,10 @@ fn emit_directives<'input>(
let f_body = match &f.body {
Some(f) => f,
None => {
- if f.import_as.is_some() {
- &empty_body
- } else {
+ if f.linkage == ast::LinkingDirective::NONE {
continue;
+ } else {
+ &empty_body
}
}
};
@@ -607,33 +607,38 @@ fn emit_directives<'input>(
}
}
emit_function_body_ops(builder, map, id_defs, opencl_id, &f_body)?;
+ emit_function_linkage(builder, id_defs, f, fn_id);
builder.select_block(None)?;
builder.end_function()?;
- if let (
- ast::MethodDeclaration {
- name: ast::MethodName::Func(fn_id),
- ..
- },
- Some(name),
- ) = (&*func_decl, &f.import_as)
- {
- builder.decorate(
- *fn_id,
- spirv::Decoration::LinkageAttributes,
- [
- dr::Operand::LiteralString(name.clone()),
- dr::Operand::LinkageType(spirv::LinkageType::Import),
- ]
- .iter()
- .cloned(),
- );
- }
}
}
}
Ok(())
}
+fn emit_function_linkage<'input>(
+ builder: &mut dr::Builder,
+ id_defs: &GlobalStringIdResolver<'input>,
+ f: &Function,
+ fn_name: spirv::Word,
+) -> Result<(), TranslateError> {
+ if f.linkage == ast::LinkingDirective::NONE {
+ return Ok(());
+ };
+ let linking_name = f.import_as.as_deref().map_or_else(
+ || match f.func_decl.borrow().name {
+ ast::MethodName::Kernel(kernel_name) => Ok(kernel_name),
+ ast::MethodName::Func(fn_id) => match id_defs.reverse_variables.get(&fn_id) {
+ Some(fn_name) => Ok(fn_name),
+ None => Err(error_unknown_symbol()),
+ },
+ },
+ Result::Ok,
+ )?;
+ emit_linking_decoration(builder, id_defs, Some(linking_name), fn_name, f.linkage);
+ Ok(())
+}
+
fn get_kernels_call_map<'input>(
module: &[Directive<'input>],
) -> HashMap<&'input str, HashSet<spirv::Word>> {
@@ -763,6 +768,7 @@ fn convert_dynamic_shared_memory_usage<'input>(
body: Some(statements),
import_as,
tuning,
+ linkage,
}) => {
let call_key = (*func_decl).borrow().name;
let statements = statements
@@ -786,6 +792,7 @@ fn convert_dynamic_shared_memory_usage<'input>(
body: Some(statements),
import_as,
tuning,
+ linkage,
})
}
directive => directive,
@@ -804,6 +811,7 @@ fn convert_dynamic_shared_memory_usage<'input>(
body: Some(statements),
import_as,
tuning,
+ linkage,
}) => {
if !methods_using_extern_shared.contains(&(*func_decl).borrow().name) {
return Directive::Method(Function {
@@ -812,6 +820,7 @@ fn convert_dynamic_shared_memory_usage<'input>(
body: Some(statements),
import_as,
tuning,
+ linkage,
});
}
let shared_id_param = new_id();
@@ -832,6 +841,7 @@ fn convert_dynamic_shared_memory_usage<'input>(
body: Some(statements),
import_as,
tuning,
+ linkage,
})
}
directive => directive,
@@ -1150,8 +1160,8 @@ fn translate_directive<'input, 'a>(
array_init: var.array_init,
},
)),
- ast::Directive::Method(_, f) => {
- translate_function(id_defs, ptx_impl_imports, f)?.map(Directive::Method)
+ ast::Directive::Method(linkage, f) => {
+ translate_function(id_defs, ptx_impl_imports, linkage, f)?.map(Directive::Method)
}
})
}
@@ -1159,6 +1169,7 @@ fn translate_directive<'input, 'a>(
fn translate_function<'input, 'a>(
id_defs: &'a mut GlobalStringIdResolver<'input>,
ptx_impl_imports: &'a mut HashMap<String, Directive<'input>>,
+ linkage: ast::LinkingDirective,
f: ast::ParsedFunction<'input>,
) -> Result<Option<Function<'input>>, TranslateError> {
let import_as = match &f.func_directive {
@@ -1178,6 +1189,7 @@ fn translate_function<'input, 'a>(
fn_decl,
f.body,
f.tuning,
+ linkage,
)?;
func.import_as = import_as;
if func.import_as.is_some() {
@@ -1213,6 +1225,7 @@ fn to_ssa<'input, 'b>(
func_decl: Rc<RefCell<ast::MethodDeclaration<'input, spirv::Word>>>,
f_body: Option<Vec<ast::Statement<ast::ParsedArgParams<'input>>>>,
tuning: Vec<ast::TuningDirective>,
+ linkage: ast::LinkingDirective,
) -> Result<Function<'input>, TranslateError> {
//deparamize_function_decl(&func_decl)?;
let f_body = match f_body {
@@ -1224,6 +1237,7 @@ fn to_ssa<'input, 'b>(
globals: Vec::new(),
import_as: None,
tuning,
+ linkage,
})
}
};
@@ -1255,6 +1269,7 @@ fn to_ssa<'input, 'b>(
body: Some(f_body),
import_as: None,
tuning,
+ linkage,
})
}
@@ -1832,6 +1847,7 @@ fn register_external_fn_call<'a>(
body: None,
import_as: Some(entry.key().clone()),
tuning: Vec::new(),
+ linkage: ast::LinkingDirective::EXTERN,
};
entry.insert(Directive::Method(func));
Ok(fn_id)
@@ -3464,37 +3480,40 @@ fn emit_variable<'input>(
[dr::Operand::LiteralInt32(align)].iter().cloned(),
);
}
- emit_linking_decoration(builder, id_defs, var.name, linking);
+ emit_linking_decoration(builder, id_defs, None, var.name, linking);
Ok(())
}
fn emit_linking_decoration<'input>(
builder: &mut dr::Builder,
id_defs: &GlobalStringIdResolver<'input>,
+ name_override: Option<&str>,
name: spirv::Word,
linking: ast::LinkingDirective,
) {
- if linking.contains(ast::LinkingDirective::EXTERN) {
- let external_name = id_defs.reverse_variables.get(&name).unwrap();
+ if linking == ast::LinkingDirective::NONE {
+ return;
+ }
+ let string_name =
+ name_override.unwrap_or_else(|| id_defs.reverse_variables.get(&name).unwrap());
+ if linking.contains(ast::LinkingDirective::VISIBLE) {
builder.decorate(
name,
spirv::Decoration::LinkageAttributes,
[
- dr::Operand::LiteralString(external_name.to_string()),
- dr::Operand::LinkageType(spirv::LinkageType::Import),
+ dr::Operand::LiteralString(string_name.to_string()),
+ dr::Operand::LinkageType(spirv::LinkageType::Export),
]
.iter()
.cloned(),
);
- }
- if linking.contains(ast::LinkingDirective::VISIBLE) {
- let external_name = id_defs.reverse_variables.get(&name).unwrap();
+ } else if linking.contains(ast::LinkingDirective::EXTERN) {
builder.decorate(
name,
spirv::Decoration::LinkageAttributes,
[
- dr::Operand::LiteralString(external_name.to_string()),
- dr::Operand::LinkageType(spirv::LinkageType::Export),
+ dr::Operand::LiteralString(string_name.to_string()),
+ dr::Operand::LinkageType(spirv::LinkageType::Import),
]
.iter()
.cloned(),
@@ -5774,6 +5793,7 @@ struct Function<'input> {
pub body: Option<Vec<ExpandedStatement>>,
import_as: Option<String>,
tuning: Vec<ast::TuningDirective>,
+ linkage: ast::LinkingDirective,
}
pub trait ArgumentMapVisitor<T: ArgParamsEx, U: ArgParamsEx> {
generated by cgit v1.2.3 (git 2.39.1) at 2024-12-25 14:25:24 +0000