diff --git a/hip_runtime-sys/build.rs b/hip_runtime-sys/build.rs index af3f787..9c3d9b4 100644 --- a/hip_runtime-sys/build.rs +++ b/hip_runtime-sys/build.rs @@ -4,5 +4,6 @@ fn main() -> Result<(), VarError> { println!("cargo:rustc-link-lib=dylib=amdhip64"); //println!("cargo:rustc-link-search=native=/opt/rocm/lib/"); println!("cargo:rustc-link-search=native=/home/ubuntu/hipamd/build/lib"); + println!("cargo:rustc-link-search=native=/home/vosen/hipamd/build/lib"); Ok(()) } diff --git a/ptx/src/ptx.lalrpop b/ptx/src/ptx.lalrpop index b20a30a..abefdf8 100644 --- a/ptx/src/ptx.lalrpop +++ b/ptx/src/ptx.lalrpop @@ -607,41 +607,33 @@ SharedVariable: ast::Variable<&'input str> = { } ModuleVariable: (ast::LinkingDirective, ast::Variable<&'input str>) = { - ".global" => { + => { let (align, v_type, name, array_init) = def; - let state_space = ast::StateSpace::Global; (linking, ast::Variable { align, v_type, state_space, name, array_init }) }, - ".shared" => { - let (align, v_type, name, array_init) = def; - let state_space = ast::StateSpace::Shared; - (linking, ast::Variable { align, v_type, state_space, name, array_init: Vec::new() }) - }, - > > =>? { + > =>? { let (align, t, name, arr_or_ptr) = var; let (v_type, state_space, array_init) = match arr_or_ptr { ast::ArrayOrPointer::Array { dimensions, init } => { - if space == ".global" { - (ast::Type::Array(t, dimensions), ast::StateSpace::Global, init) - } else { - (ast::Type::Array(t, dimensions), ast::StateSpace::Shared, init) - } + (ast::Type::Array(t, dimensions), space, init) } ast::ArrayOrPointer::Pointer => { if !linking.contains(ast::LinkingDirective::EXTERN) { return Err(ParseError::User { error: ast::PtxError::NonExternPointer }); } - if space == ".global" { - (ast::Type::Array(t, Vec::new()), ast::StateSpace::Global, Vec::new()) - } else { - (ast::Type::Array(t, Vec::new()), ast::StateSpace::Shared, Vec::new()) - } + (ast::Type::Array(t, Vec::new()), space, Vec::new()) } }; Ok((linking, ast::Variable{ align, v_type, state_space, name, array_init })) } } +VariableStateSpace: ast::StateSpace = { + ".const" => ast::StateSpace::Const, + ".global" => ast::StateSpace::Global, + ".shared" => ast::StateSpace::Shared, +}; + // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parameter-state-space ParamVariable: (Option, Vec, ast::Type, &'input str) = { ".param" > => { @@ -2096,4 +2088,11 @@ CommaNonEmpty: Vec = { Or: T1 = { T1, T2 +} + +#[inline] +Or3: T1 = { + T1, + T2, + T3 } \ No newline at end of file diff --git a/ptx/src/test/spirv_run/const.ptx b/ptx/src/test/spirv_run/const.ptx new file mode 100644 index 0000000..c22ac2b --- /dev/null +++ b/ptx/src/test/spirv_run/const.ptx @@ -0,0 +1,31 @@ +.version 6.5 +.target sm_30 +.address_size 64 + +.const .align 8 .b16 constparams[4] = { 10, 20, 30, 40 }; + +.visible .entry const( + .param .u64 input, + .param .u64 output +) +{ + .reg .u64 in_addr; + .reg .u64 out_addr; + .reg .b16 temp1; + .reg .b16 temp2; + .reg .b16 temp3; + .reg .b16 temp4; + + ld.param.u64 in_addr, [input]; + ld.param.u64 out_addr, [output]; + + ld.const.b16 temp1, constparams[0]; + ld.const.b16 temp2, constparams[1]; + ld.const.b16 temp3, constparams[2]; + ld.const.b16 temp4, constparams[3]; + st.u16 [out_addr], temp1; + st.u16 [out_addr+2], temp2; + st.u16 [out_addr+4], temp3; + st.u16 [out_addr+6], temp4; + ret; +} diff --git a/ptx/src/test/spirv_run/const.spvtxt b/ptx/src/test/spirv_run/const.spvtxt new file mode 100644 index 0000000..9a7f254 --- /dev/null +++ b/ptx/src/test/spirv_run/const.spvtxt @@ -0,0 +1,47 @@ + OpCapability GenericPointer + OpCapability Linkage + OpCapability Addresses + OpCapability Kernel + OpCapability Int8 + OpCapability Int16 + OpCapability Int64 + OpCapability Float16 + OpCapability Float64 + %21 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %1 "clz" + %void = OpTypeVoid + %ulong = OpTypeInt 64 0 + %24 = OpTypeFunction %void %ulong %ulong +%_ptr_Function_ulong = OpTypePointer Function %ulong + %uint = OpTypeInt 32 0 +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Generic_uint = OpTypePointer Generic %uint + %1 = OpFunction %void None %24 + %7 = OpFunctionParameter %ulong + %8 = OpFunctionParameter %ulong + %19 = 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_uint Function + OpStore %2 %7 + OpStore %3 %8 + %9 = OpLoad %ulong %2 Aligned 8 + OpStore %4 %9 + %10 = OpLoad %ulong %3 Aligned 8 + OpStore %5 %10 + %12 = OpLoad %ulong %4 + %17 = OpConvertUToPtr %_ptr_Generic_uint %12 + %11 = OpLoad %uint %17 Aligned 4 + OpStore %6 %11 + %14 = OpLoad %uint %6 + %13 = OpExtInst %uint %21 clz %14 + OpStore %6 %13 + %15 = OpLoad %ulong %5 + %16 = OpLoad %uint %6 + %18 = OpConvertUToPtr %_ptr_Generic_uint %15 + OpStore %18 %16 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs index 51f1930..13cf0f1 100644 --- a/ptx/src/test/spirv_run/mod.rs +++ b/ptx/src/test/spirv_run/mod.rs @@ -203,6 +203,7 @@ test_ptx!( ); test_ptx!(non_scalar_ptr_offset, [1u32, 2u32, 3u32, 4u32], [7u32]); test_ptx!(stateful_neg_offset, [1237518u64], [1237518u64]); +test_ptx!(const, [0u16], [10u16, 20, 30, 40]); struct DisplayError { err: T,