Add support for blocks

ptx/src/test/spirv_run/block.ptx

@@ -0,0 +1,26 @@
+.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry block(
+	.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.u64          temp, [in_addr];
+	add.u64		    temp2, temp, 1;
+    {
+        .reg .u64 	    temp2;
+	    add.u64		    temp2, temp2, 1;
+    }
+    st.u64          [out_addr], temp2;
+	ret;
+}

ptx/src/test/spirv_run/block.spvtxt

@@ -0,0 +1,44 @@
+               OpCapability GenericPointer
+               OpCapability Linkage
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int64
+               OpCapability Int8
+          %1 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %5 "block"
+       %void = OpTypeVoid
+      %ulong = OpTypeInt 64 0
+          %4 = OpTypeFunction %void %ulong %ulong
+%_ptr_Function_ulong = OpTypePointer Function %ulong
+%_ptr_Generic_ulong = OpTypePointer Generic %ulong
+    %ulong_1 = OpConstant %ulong 1
+  %ulong_1_0 = OpConstant %ulong 1
+          %5 = OpFunction %void None %4
+          %6 = OpFunctionParameter %ulong
+          %7 = OpFunctionParameter %ulong
+         %25 = OpLabel
+          %8 = OpVariable %_ptr_Function_ulong Function
+          %9 = OpVariable %_ptr_Function_ulong Function
+         %10 = OpVariable %_ptr_Function_ulong Function
+         %11 = OpVariable %_ptr_Function_ulong Function
+               OpStore %8 %6
+               OpStore %9 %7
+         %14 = OpLoad %ulong %8
+         %23 = OpConvertUToPtr %_ptr_Generic_ulong %14
+         %13 = OpLoad %ulong %23
+               OpStore %10 %13
+         %16 = OpLoad %ulong %10
+         %15 = OpIAdd %ulong %16 %ulong_1
+               OpStore %11 %15
+         %12 = OpVariable %_ptr_Function_ulong Function
+         %18 = OpLoad %ulong %12
+         %17 = OpIAdd %ulong %18 %ulong_1_0
+               OpStore %12 %17
+         %19 = OpLoad %ulong %9
+         %20 = OpLoad %ulong %11
+         %24 = OpConvertUToPtr %_ptr_Generic_ulong %19
+               OpStore %24 %20
+               OpReturn
+               OpFunctionEnd
+               

ptx/src/test/spirv_run/mod.rs

@@ -50,6 +50,7 @@ test_ptx!(not, [0u64], [u64::max_value()]);
 test_ptx!(shl, [11u64], [44u64]);
 test_ptx!(cvt_sat_s_u, [-1i32], [0i32]);
 test_ptx!(cvta, [3.0f32], [3.0f32]);
+test_ptx!(block, [1u64], [2u64]);
 
 struct DisplayError<T: Debug> {
     err: T,

ptx/src/translate.rs

@@ -271,7 +271,6 @@ fn normalize_predicates(
     let mut result = Vec::with_capacity(func.len());
     for s in func {
         match s {
-            ast::Statement::Block(_) => todo!(),
             ast::Statement::Label(id) => result.push(Statement::Label(id)),
             ast::Statement::Instruction(pred, inst) => {
                 if let Some(pred) = pred {
@@ -302,6 +301,8 @@ fn normalize_predicates(
             ast::Statement::Variable(var) => {
                 result.push(Statement::Variable(var.name, var.v_type, var.space))
             }
+            // Blocks are flattened when resolving ids
+            ast::Statement::Block(_) => unreachable!(),
         }
     }
     result
@@ -659,7 +660,7 @@ fn emit_function_body_ops(
             }
             Statement::Instruction(inst) => match inst {
                 ast::Instruction::Abs(_, _) => todo!(),
-                ast::Instruction::Call(_,_) => todo!(),
+                ast::Instruction::Call(_, _) => todo!(),
                 // SPIR-V does not support marking jumps as guaranteed-converged
                 ast::Instruction::Bra(_, arg) => {
                     builder.branch(arg.src)?;
@@ -1084,7 +1085,13 @@ fn expand_map_variables<'a>(
     s: ast::Statement<ast::ParsedArgParams<'a>>,
 ) {
     match s {
-        ast::Statement::Block(_) => todo!(),
+        ast::Statement::Block(block) => {
+            id_defs.start_block();
+            for s in block {
+                expand_map_variables(id_defs, result, s);
+            }
+            id_defs.end_block();
+        }
         ast::Statement::Label(name) => result.push(ast::Statement::Label(id_defs.get_id(name))),
         ast::Statement::Instruction(p, i) => result.push(ast::Statement::Instruction(
             p.map(|p| p.map_variable(&mut |id| id_defs.get_id(id))),
@@ -1118,7 +1125,7 @@ fn expand_map_variables<'a>(
 
 struct StringIdResolver<'a> {
     current_id: spirv::Word,
-    variables: HashMap<Cow<'a, str>, spirv::Word>,
+    variables: Vec<HashMap<Cow<'a, str>, spirv::Word>>,
     type_check: HashMap<u32, ast::Type>,
 }
 
@@ -1126,7 +1133,7 @@ impl<'a> StringIdResolver<'a> {
     fn new() -> Self {
         StringIdResolver {
             current_id: 0u32,
-            variables: HashMap::new(),
+            variables: vec![HashMap::new(); 1],
             type_check: HashMap::new(),
         }
     }
@@ -1138,13 +1145,30 @@ impl<'a> StringIdResolver<'a> {
         }
     }
 
+    fn start_block(&mut self) {
+        self.variables.push(HashMap::new())
+    }
+
+    fn end_block(&mut self) {
+        self.variables.pop();
+    }
+
     fn get_id(&self, id: &'a str) -> spirv::Word {
-        self.variables[id]
+        for scope in self.variables.iter().rev() {
+            match scope.get(id) {
+                Some(id) => return *id,
+                None => continue,
+            }
+        }
+        panic!()
     }
 
     fn add_def(&mut self, id: &'a str, typ: Option<ast::Type>) -> spirv::Word {
         let numeric_id = self.current_id;
-        self.variables.insert(Cow::Borrowed(id), numeric_id);
+        self.variables
+            .last_mut()
+            .unwrap()
+            .insert(Cow::Borrowed(id), numeric_id);
         if let Some(typ) = typ {
             self.type_check.insert(numeric_id, typ);
         }
@@ -1162,6 +1186,8 @@ impl<'a> StringIdResolver<'a> {
         let numeric_id = self.current_id;
         for i in 0..count {
             self.variables
+                .last_mut()
+                .unwrap()
                 .insert(Cow::Owned(format!("{}{}", base_id, i)), numeric_id + i);
             self.type_check.insert(numeric_id + i, typ);
         }