Parse and test const buffers

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(())
 }

ptx/src/ptx.lalrpop

@@ -607,41 +607,33 @@ SharedVariable: ast::Variable<&'input str> = {
 }
 
 ModuleVariable: (ast::LinkingDirective, ast::Variable<&'input str>) = {
-    <linking:LinkingDirectives> ".global" <def:GlobalVariableDefinitionNoArray> => {
+    <linking:LinkingDirectives> <state_space:VariableStateSpace> <def:GlobalVariableDefinitionNoArray> => {
         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 })
     },
-    <linking:LinkingDirectives> ".shared" <def:GlobalVariableDefinitionNoArray> => {
+    <linking:LinkingDirectives> <space:VariableStateSpace> <var:VariableArrayOrPointer<SizedScalarType>> =>? {
-        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() })
-    },
-    <linking:LinkingDirectives> <space:Or<".global", ".shared">> <var:VariableArrayOrPointer<SizedScalarType>> =>? {
         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), space, init)
-                    (ast::Type::Array(t, dimensions), ast::StateSpace::Global, init)
-                } else {
-                    (ast::Type::Array(t, dimensions), ast::StateSpace::Shared, 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()), space, Vec::new())
-                    (ast::Type::Array(t, Vec::new()), ast::StateSpace::Global, Vec::new())
-                } else {
-                    (ast::Type::Array(t, Vec::new()), ast::StateSpace::Shared, 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<u32>, Vec<u8>, ast::Type, &'input str) = {
     ".param" <var:VariableScalar<LdStScalarType>> => {
@@ -2096,4 +2088,11 @@ CommaNonEmpty<T>: Vec<T> = {
 Or<T1, T2>: T1 = {
     T1,
     T2
+}
+
+#[inline]
+Or3<T1, T2, T3>: T1 = {
+    T1,
+    T2,
+    T3
 }

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

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

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<T: Debug> {
     err: T,