Fix linkage

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