From afe912086895bb6a9c5d7f386fd8e7dd90561d63 Mon Sep 17 00:00:00 2001 From: Andrzej Janik Date: Sat, 18 Sep 2021 22:49:00 +0000 Subject: Fix linkage --- ptx/src/test/spirv_run/add_tuning.spvtxt | 103 ++++++++++++++++-------------- ptx/src/test/spirv_run/extern_func.ptx | 39 +++++++++++ ptx/src/test/spirv_run/extern_func.spvtxt | 75 ++++++++++++++++++++++ ptx/src/test/spirv_run/func_ptr.spvtxt | 60 +++++++++-------- ptx/src/test/spirv_run/mod.rs | 17 ++++- ptx/src/translate.rs | 88 +++++++++++++++---------- 6 files changed, 270 insertions(+), 112 deletions(-) create mode 100644 ptx/src/test/spirv_run/extern_func.ptx create mode 100644 ptx/src/test/spirv_run/extern_func.spvtxt 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> { + 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 { 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> { @@ -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>, + linkage: ast::LinkingDirective, f: ast::ParsedFunction<'input>, ) -> Result>, 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>>, f_body: Option>>>, tuning: Vec, + linkage: ast::LinkingDirective, ) -> Result, 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>, import_as: Option, tuning: Vec, + linkage: ast::LinkingDirective, } pub trait ArgumentMapVisitor { -- cgit v1.2.3