Archive/ZLUDA
0
0
Fork 0
mirror of https://github.com/vosen/ZLUDA.git synced 2025-04-28 21:47:57 +03:00
Code Issues Projects Releases Packages Wiki Activity Actions

Fix linkage

Browse source
This commit is contained in:
Andrzej Janik 2021-09-18 22:49:00 +00:00
parent 04a411fe22
commit afe9120868
6 changed files with 270 additions and 112 deletions
Download patch file Download diff file Expand all files Collapse all files

103
ptx/src/test/spirv_run/add_tuning.spvtxt
View file

@ -1,48 +1,55 @@
OpCapability GenericPointer ; SPIR-V
OpCapability Linkage ; Version: 1.3
OpCapability Addresses ; Generator: rspirv
OpCapability Kernel ; Bound: 29
OpCapability Int8 OpCapability GenericPointer
OpCapability Int16 OpCapability Linkage
OpCapability Int64 OpCapability Addresses
OpCapability Float16 OpCapability Kernel
OpCapability Float64 OpCapability Int8
%23 = OpExtInstImport "OpenCL.std" OpCapability Int16
OpMemoryModel Physical64 OpenCL OpCapability Int64
OpEntryPoint Kernel %1 "add_tuning" OpCapability Float16
OpExecutionMode %1 MaxWorkgroupSizeINTEL 256 1 1 OpCapability Float64
%void = OpTypeVoid OpCapability DenormFlushToZero
%ulong = OpTypeInt 64 0 %23 = OpExtInstImport "OpenCL.std"
%26 = OpTypeFunction %void %ulong %ulong OpMemoryModel Physical64 OpenCL
%_ptr_Function_ulong = OpTypePointer Function %ulong OpEntryPoint Kernel %1 "add_tuning"
%_ptr_Generic_ulong = OpTypePointer Generic %ulong OpExecutionMode %1 ContractionOff
%ulong_1 = OpConstant %ulong 1 ; OpExecutionMode %1 MaxWorkgroupSizeINTEL 256 1 1
%1 = OpFunction %void None %26 OpDecorate %1 LinkageAttributes "add_tuning" Export
%8 = OpFunctionParameter %ulong %24 = OpTypeVoid
%9 = OpFunctionParameter %ulong %25 = OpTypeInt 64 0
%21 = OpLabel %26 = OpTypeFunction %24 %25 %25
%2 = OpVariable %_ptr_Function_ulong Function %27 = OpTypePointer Function %25
%3 = OpVariable %_ptr_Function_ulong Function %28 = OpTypePointer Generic %25
%4 = OpVariable %_ptr_Function_ulong Function %18 = OpConstant %25 1
%5 = OpVariable %_ptr_Function_ulong Function %1 = OpFunction %24 None %26
%6 = OpVariable %_ptr_Function_ulong Function %8 = OpFunctionParameter %25
%7 = OpVariable %_ptr_Function_ulong Function %9 = OpFunctionParameter %25
OpStore %2 %8 %21 = OpLabel
OpStore %3 %9 %2 = OpVariable %27 Function
%10 = OpLoad %ulong %2 Aligned 8 %3 = OpVariable %27 Function
OpStore %4 %10 %4 = OpVariable %27 Function
%11 = OpLoad %ulong %3 Aligned 8 %5 = OpVariable %27 Function
OpStore %5 %11 %6 = OpVariable %27 Function
%13 = OpLoad %ulong %4 %7 = OpVariable %27 Function
%19 = OpConvertUToPtr %_ptr_Generic_ulong %13 OpStore %2 %8
%12 = OpLoad %ulong %19 Aligned 8 OpStore %3 %9
OpStore %6 %12 %10 = OpLoad %25 %2 Aligned 8
%15 = OpLoad %ulong %6 OpStore %4 %10
%14 = OpIAdd %ulong %15 %ulong_1 %11 = OpLoad %25 %3 Aligned 8
OpStore %7 %14 OpStore %5 %11
%16 = OpLoad %ulong %5 %13 = OpLoad %25 %4
%17 = OpLoad %ulong %7 %19 = OpConvertUToPtr %28 %13
%20 = OpConvertUToPtr %_ptr_Generic_ulong %16 %12 = OpLoad %25 %19 Aligned 8
OpStore %20 %17 Aligned 8 OpStore %6 %12
OpReturn %15 = OpLoad %25 %6
OpFunctionEnd %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

39
ptx/src/test/spirv_run/extern_func.ptx Normal file
View file

@ -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;
}

75
ptx/src/test/spirv_run/extern_func.spvtxt Normal file
View file

@ -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

60
ptx/src/test/spirv_run/func_ptr.spvtxt
View file

@ -7,24 +7,26 @@
OpCapability Int64 OpCapability Int64
OpCapability Float16 OpCapability Float16
OpCapability Float64 OpCapability Float64
%38 = OpExtInstImport "OpenCL.std" OpCapability DenormFlushToZero
%39 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %11 "func_ptr" OpEntryPoint Kernel %12 "func_ptr"
OpExecutionMode %11 ContractionOff OpExecutionMode %12 ContractionOff
OpDecorate %12 LinkageAttributes "func_ptr" Export
%void = OpTypeVoid %void = OpTypeVoid
%float = OpTypeFloat 32 %float = OpTypeFloat 32
%41 = OpTypeFunction %float %float %float %42 = OpTypeFunction %float %float %float
%_ptr_Function_float = OpTypePointer Function %float %_ptr_Function_float = OpTypePointer Function %float
%ulong = OpTypeInt 64 0 %ulong = OpTypeInt 64 0
%44 = OpTypeFunction %void %ulong %ulong %45 = OpTypeFunction %void %ulong %ulong
%_ptr_Function_ulong = OpTypePointer Function %ulong %_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Generic_ulong = OpTypePointer Generic %ulong %_ptr_Generic_ulong = OpTypePointer Generic %ulong
%ulong_1 = OpConstant %ulong 1 %ulong_1 = OpConstant %ulong 1
%ulong_0 = OpConstant %ulong 0 %ulong_0 = OpConstant %ulong 0
%1 = OpFunction %float None %41 %1 = OpFunction %float None %42
%5 = OpFunctionParameter %float %5 = OpFunctionParameter %float
%6 = OpFunctionParameter %float %6 = OpFunctionParameter %float
%10 = OpLabel %11 = OpLabel
%3 = OpVariable %_ptr_Function_float Function %3 = OpVariable %_ptr_Function_float Function
%4 = OpVariable %_ptr_Function_float Function %4 = OpVariable %_ptr_Function_float Function
%2 = OpVariable %_ptr_Function_float Function %2 = OpVariable %_ptr_Function_float Function
@ -34,40 +36,42 @@
%9 = OpLoad %float %4 %9 = OpLoad %float %4
%7 = OpFAdd %float %8 %9 %7 = OpFAdd %float %8 %9
OpStore %2 %7 OpStore %2 %7
%10 = OpLoad %float %2
OpReturnValue %10
OpFunctionEnd OpFunctionEnd
%11 = OpFunction %void None %44 %12 = OpFunction %void None %45
%19 = OpFunctionParameter %ulong
%20 = OpFunctionParameter %ulong %20 = OpFunctionParameter %ulong
%36 = OpLabel %21 = OpFunctionParameter %ulong
%12 = OpVariable %_ptr_Function_ulong Function %37 = OpLabel
%13 = OpVariable %_ptr_Function_ulong Function %13 = OpVariable %_ptr_Function_ulong Function
%14 = OpVariable %_ptr_Function_ulong Function %14 = OpVariable %_ptr_Function_ulong Function
%15 = OpVariable %_ptr_Function_ulong Function %15 = OpVariable %_ptr_Function_ulong Function
%16 = OpVariable %_ptr_Function_ulong Function %16 = OpVariable %_ptr_Function_ulong Function
%17 = OpVariable %_ptr_Function_ulong Function %17 = OpVariable %_ptr_Function_ulong Function
%18 = OpVariable %_ptr_Function_ulong Function %18 = OpVariable %_ptr_Function_ulong Function
OpStore %12 %19 %19 = OpVariable %_ptr_Function_ulong Function
OpStore %13 %20 OpStore %13 %20
%21 = OpLoad %ulong %12 Aligned 8
OpStore %14 %21 OpStore %14 %21
%22 = OpLoad %ulong %13 Aligned 8 %22 = OpLoad %ulong %13 Aligned 8
OpStore %15 %22 OpStore %15 %22
%24 = OpLoad %ulong %14 %23 = OpLoad %ulong %14 Aligned 8
%34 = OpConvertUToPtr %_ptr_Generic_ulong %24
%23 = OpLoad %ulong %34 Aligned 8
OpStore %16 %23 OpStore %16 %23
%26 = OpLoad %ulong %16 %25 = OpLoad %ulong %15
%25 = OpIAdd %ulong %26 %ulong_1 %35 = OpConvertUToPtr %_ptr_Generic_ulong %25
OpStore %17 %25 %24 = OpLoad %ulong %35 Aligned 8
%27 = OpCopyObject %ulong %ulong_0 OpStore %17 %24
OpStore %18 %27 %27 = OpLoad %ulong %17
%29 = OpLoad %ulong %17 %26 = OpIAdd %ulong %27 %ulong_1
OpStore %18 %26
%28 = OpCopyObject %ulong %ulong_0
OpStore %19 %28
%30 = OpLoad %ulong %18 %30 = OpLoad %ulong %18
%28 = OpIAdd %ulong %29 %30 %31 = OpLoad %ulong %19
OpStore %17 %28 %29 = OpIAdd %ulong %30 %31
%31 = OpLoad %ulong %15 OpStore %18 %29
%32 = OpLoad %ulong %17 %32 = OpLoad %ulong %16
%35 = OpConvertUToPtr %_ptr_Generic_ulong %31 %33 = OpLoad %ulong %18
OpStore %35 %32 Aligned 8 %36 = OpConvertUToPtr %_ptr_Generic_ulong %32
OpStore %36 %33 Aligned 8
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

17
ptx/src/test/spirv_run/mod.rs
View file

@ -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]); 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!(prmt, [0x70c507d6u32, 0x6fbd4b5cu32], [0x6fbdd65cu32]);
test_ptx!(activemask, [0u32], [1u32]); test_ptx!(activemask, [0u32], [1u32]);
test_ptx!(membar, [152731u32], [152731u32]); test_ptx!(membar, [152731u32], [152731u32]);
test_ptx!(func_ptr, [152731u64], [152732u64]); test_ptx!(func_ptr);
test_ptx!(lanemask_lt, [187235u32], [187236u32]); test_ptx!(lanemask_lt);
test_ptx!(extern_func);
struct DisplayError<T: Debug> { struct DisplayError<T: Debug> {
err: T, err: T,

88
ptx/src/translate.rs
View file

@ -543,10 +543,10 @@ fn emit_directives<'input>(
let f_body = match &f.body { let f_body = match &f.body {
Some(f) => f, Some(f) => f,
None => { None => {
if f.import_as.is_some() { if f.linkage == ast::LinkingDirective::NONE {
&empty_body
} else {
continue; 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_body_ops(builder, map, id_defs, opencl_id, &f_body)?;
emit_function_linkage(builder, id_defs, f, fn_id);
builder.select_block(None)?; builder.select_block(None)?;
builder.end_function()?; 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(()) 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>( fn get_kernels_call_map<'input>(
module: &[Directive<'input>], module: &[Directive<'input>],
) -> HashMap<&'input str, HashSet<spirv::Word>> { ) -> HashMap<&'input str, HashSet<spirv::Word>> {
@ -763,6 +768,7 @@ fn convert_dynamic_shared_memory_usage<'input>(
body: Some(statements), body: Some(statements),
import_as, import_as,
tuning, tuning,
linkage,
}) => { }) => {
let call_key = (*func_decl).borrow().name; let call_key = (*func_decl).borrow().name;
let statements = statements let statements = statements
@ -786,6 +792,7 @@ fn convert_dynamic_shared_memory_usage<'input>(
body: Some(statements), body: Some(statements),
import_as, import_as,
tuning, tuning,
linkage,
}) })
} }
directive => directive, directive => directive,
@ -804,6 +811,7 @@ fn convert_dynamic_shared_memory_usage<'input>(
body: Some(statements), body: Some(statements),
import_as, import_as,
tuning, tuning,
linkage,
}) => { }) => {
if !methods_using_extern_shared.contains(&(*func_decl).borrow().name) { if !methods_using_extern_shared.contains(&(*func_decl).borrow().name) {
return Directive::Method(Function { return Directive::Method(Function {
@ -812,6 +820,7 @@ fn convert_dynamic_shared_memory_usage<'input>(
body: Some(statements), body: Some(statements),
import_as, import_as,
tuning, tuning,
linkage,
}); });
} }
let shared_id_param = new_id(); let shared_id_param = new_id();
@ -832,6 +841,7 @@ fn convert_dynamic_shared_memory_usage<'input>(
body: Some(statements), body: Some(statements),
import_as, import_as,
tuning, tuning,
linkage,
}) })
} }
directive => directive, directive => directive,
@ -1150,8 +1160,8 @@ fn translate_directive<'input, 'a>(
array_init: var.array_init, array_init: var.array_init,
}, },
)), )),
ast::Directive::Method(_, f) => { ast::Directive::Method(linkage, f) => {
translate_function(id_defs, ptx_impl_imports, f)?.map(Directive::Method) 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>( fn translate_function<'input, 'a>(
id_defs: &'a mut GlobalStringIdResolver<'input>, id_defs: &'a mut GlobalStringIdResolver<'input>,
ptx_impl_imports: &'a mut HashMap<String, Directive<'input>>, ptx_impl_imports: &'a mut HashMap<String, Directive<'input>>,
linkage: ast::LinkingDirective,
f: ast::ParsedFunction<'input>, f: ast::ParsedFunction<'input>,
) -> Result<Option<Function<'input>>, TranslateError> { ) -> Result<Option<Function<'input>>, TranslateError> {
let import_as = match &f.func_directive { let import_as = match &f.func_directive {
@ -1178,6 +1189,7 @@ fn translate_function<'input, 'a>(
fn_decl, fn_decl,
f.body, f.body,
f.tuning, f.tuning,
linkage,
)?; )?;
func.import_as = import_as; func.import_as = import_as;
if func.import_as.is_some() { if func.import_as.is_some() {
@ -1213,6 +1225,7 @@ fn to_ssa<'input, 'b>(
func_decl: Rc<RefCell<ast::MethodDeclaration<'input, spirv::Word>>>, func_decl: Rc<RefCell<ast::MethodDeclaration<'input, spirv::Word>>>,
f_body: Option<Vec<ast::Statement<ast::ParsedArgParams<'input>>>>, f_body: Option<Vec<ast::Statement<ast::ParsedArgParams<'input>>>>,
tuning: Vec<ast::TuningDirective>, tuning: Vec<ast::TuningDirective>,
linkage: ast::LinkingDirective,
) -> Result<Function<'input>, TranslateError> { ) -> Result<Function<'input>, TranslateError> {
//deparamize_function_decl(&func_decl)?; //deparamize_function_decl(&func_decl)?;
let f_body = match f_body { let f_body = match f_body {
@ -1224,6 +1237,7 @@ fn to_ssa<'input, 'b>(
globals: Vec::new(), globals: Vec::new(),
import_as: None, import_as: None,
tuning, tuning,
linkage,
}) })
} }
}; };
@ -1255,6 +1269,7 @@ fn to_ssa<'input, 'b>(
body: Some(f_body), body: Some(f_body),
import_as: None, import_as: None,
tuning, tuning,
linkage,
}) })
} }
@ -1832,6 +1847,7 @@ fn register_external_fn_call<'a>(
body: None, body: None,
import_as: Some(entry.key().clone()), import_as: Some(entry.key().clone()),
tuning: Vec::new(), tuning: Vec::new(),
linkage: ast::LinkingDirective::EXTERN,
}; };
entry.insert(Directive::Method(func)); entry.insert(Directive::Method(func));
Ok(fn_id) Ok(fn_id)
@ -3464,37 +3480,40 @@ fn emit_variable<'input>(
[dr::Operand::LiteralInt32(align)].iter().cloned(), [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(()) Ok(())
} }
fn emit_linking_decoration<'input>( fn emit_linking_decoration<'input>(
builder: &mut dr::Builder, builder: &mut dr::Builder,
id_defs: &GlobalStringIdResolver<'input>, id_defs: &GlobalStringIdResolver<'input>,
name_override: Option<&str>,
name: spirv::Word, name: spirv::Word,
linking: ast::LinkingDirective, linking: ast::LinkingDirective,
) { ) {
if linking.contains(ast::LinkingDirective::EXTERN) { if linking == ast::LinkingDirective::NONE {
let external_name = id_defs.reverse_variables.get(&name).unwrap(); return;
}
let string_name =
name_override.unwrap_or_else(|| id_defs.reverse_variables.get(&name).unwrap());
if linking.contains(ast::LinkingDirective::VISIBLE) {
builder.decorate( builder.decorate(
name, name,
spirv::Decoration::LinkageAttributes, spirv::Decoration::LinkageAttributes,
[ [
dr::Operand::LiteralString(external_name.to_string()), dr::Operand::LiteralString(string_name.to_string()),
dr::Operand::LinkageType(spirv::LinkageType::Import), dr::Operand::LinkageType(spirv::LinkageType::Export),
] ]
.iter() .iter()
.cloned(), .cloned(),
); );
} } else if linking.contains(ast::LinkingDirective::EXTERN) {
if linking.contains(ast::LinkingDirective::VISIBLE) {
let external_name = id_defs.reverse_variables.get(&name).unwrap();
builder.decorate( builder.decorate(
name, name,
spirv::Decoration::LinkageAttributes, spirv::Decoration::LinkageAttributes,
[ [
dr::Operand::LiteralString(external_name.to_string()), dr::Operand::LiteralString(string_name.to_string()),
dr::Operand::LinkageType(spirv::LinkageType::Export), dr::Operand::LinkageType(spirv::LinkageType::Import),
] ]
.iter() .iter()
.cloned(), .cloned(),
@ -5774,6 +5793,7 @@ struct Function<'input> {
pub body: Option<Vec<ExpandedStatement>>, pub body: Option<Vec<ExpandedStatement>>,
import_as: Option<String>, import_as: Option<String>,
tuning: Vec<ast::TuningDirective>, tuning: Vec<ast::TuningDirective>,
linkage: ast::LinkingDirective,
} }
pub trait ArgumentMapVisitor<T: ArgParamsEx, U: ArgParamsEx> { pub trait ArgumentMapVisitor<T: ArgParamsEx, U: ArgParamsEx> {