From e7f10afb514f013f40a0b415705d3ec008454420 Mon Sep 17 00:00:00 2001 From: Violet Date: Fri, 5 Sep 2025 16:41:41 -0700 Subject: [PATCH] Use `Vec` as const/global variable initializer (#490) --- ptx/src/pass/insert_implicit_conversions2.rs | 6 +- ptx/src/pass/llvm/emit.rs | 117 ++++++++----------- ptx/src/pass/mod.rs | 2 +- ptx/src/pass/normalize_identifiers2.rs | 19 ++- ptx/src/test/ll/const_ident.ll | 55 +++++++++ ptx/src/test/ll/global_array_f32.ll | 28 +++++ ptx/src/test/spirv_run/const_ident.ptx | 39 +++++++ ptx/src/test/spirv_run/global_array_f32.ptx | 22 ++++ ptx/src/test/spirv_run/mod.rs | 2 + ptx_parser/src/ast.rs | 9 +- ptx_parser/src/lib.rs | 104 ++++++----------- 11 files changed, 255 insertions(+), 148 deletions(-) create mode 100644 ptx/src/test/ll/const_ident.ll create mode 100644 ptx/src/test/ll/global_array_f32.ll create mode 100644 ptx/src/test/spirv_run/const_ident.ptx create mode 100644 ptx/src/test/spirv_run/global_array_f32.ptx diff --git a/ptx/src/pass/insert_implicit_conversions2.rs b/ptx/src/pass/insert_implicit_conversions2.rs index 5b0fd3b..b2d3161 100644 --- a/ptx/src/pass/insert_implicit_conversions2.rs +++ b/ptx/src/pass/insert_implicit_conversions2.rs @@ -208,7 +208,7 @@ fn default_implicit_conversion_type( if should_bitcast(instruction_type, operand_type) { Ok(Some(ConversionKind::Default)) } else { - Err(TranslateError::MismatchedType) + Err(error_mismatched_type()) } } else { Ok(Some(ConversionKind::PtrToPtr)) @@ -264,14 +264,14 @@ pub(crate) fn should_convert_relaxed_dst_wrapper( (instruction_space, instruction_type): (ast::StateSpace, &ast::Type), ) -> Result, TranslateError> { if operand_space != instruction_space { - return Err(TranslateError::MismatchedType); + return Err(error_mismatched_type()); } if operand_type == instruction_type { return Ok(None); } match should_convert_relaxed_dst(operand_type, instruction_type) { conv @ Some(_) => Ok(conv), - None => Err(TranslateError::MismatchedType), + None => Err(error_mismatched_type()), } } diff --git a/ptx/src/pass/llvm/emit.rs b/ptx/src/pass/llvm/emit.rs index dc9befa..d3b81cd 100644 --- a/ptx/src/pass/llvm/emit.rs +++ b/ptx/src/pass/llvm/emit.rs @@ -24,8 +24,6 @@ // shows it fails inside amdgpu-isel. You can get a little bit furthr with "-mllvm -global-isel", // but it will too fail similarly, but with "unable to legalize instruction" -use std::array::TryFromSliceError; -use std::convert::TryInto; use std::ffi::{CStr, NulError}; use std::{i8, ptr, u64}; @@ -249,79 +247,51 @@ impl<'a, 'input> ModuleEmitContext<'a, 'input> { unsafe { LLVMSetAlignment(global, align) }; } if !var.array_init.is_empty() { - self.emit_array_init(&var.v_type, &*var.array_init, global)?; + let initializer = self.get_array_init(&var.v_type, &*var.array_init)?; + unsafe { LLVMSetInitializer(global, initializer) }; } Ok(()) } - // TODO: instead of Vec we should emit a typed initializer - fn emit_array_init( - &mut self, + fn get_array_init( + &self, type_: &ast::Type, - array_init: &[u8], - global: *mut llvm_zluda::LLVMValue, - ) -> Result<(), TranslateError> { - match type_ { + array_init: &[ast::RegOrImmediate], + ) -> Result<*mut LLVMValue, TranslateError> { + let initializer = match type_ { ast::Type::Array(None, scalar, dimensions) => { if dimensions.len() != 1 { todo!() } - if dimensions[0] as usize * scalar.size_of() as usize != array_init.len() { + if dimensions[0] as usize != array_init.len() { return Err(error_unreachable()); } let type_ = get_scalar_type(self.context, *scalar); let mut elements = array_init - .chunks(scalar.size_of() as usize) - .map(|chunk| self.constant_from_bytes(*scalar, chunk, type_)) - .collect::, _>>() - .map_err(|_| error_unreachable())?; - let initializer = - unsafe { LLVMConstArray2(type_, elements.as_mut_ptr(), elements.len() as u64) }; - unsafe { LLVMSetInitializer(global, initializer) }; + .iter() + .map(|elem| match elem { + ast::RegOrImmediate::Reg(reg) => { + Ok(unsafe { LLVMConstPtrToInt(self.resolver.value(*reg)?, type_) }) + } + ast::RegOrImmediate::Imm(imm) => { + Ok(get_immediate_value(self.context, scalar, imm)) + } + }) + .collect::, _>>()?; + unsafe { LLVMConstArray2(type_, elements.as_mut_ptr(), elements.len() as u64) } } - _ => todo!(), - } - Ok(()) - } - - fn constant_from_bytes( - &self, - scalar: ast::ScalarType, - bytes: &[u8], - llvm_type: LLVMTypeRef, - ) -> Result { - Ok(match scalar { - ptx_parser::ScalarType::Pred - | ptx_parser::ScalarType::S8 - | ptx_parser::ScalarType::B8 - | ptx_parser::ScalarType::U8 => unsafe { - LLVMConstInt(llvm_type, u8::from_le_bytes(bytes.try_into()?) as u64, 0) - }, - ptx_parser::ScalarType::S16 - | ptx_parser::ScalarType::B16 - | ptx_parser::ScalarType::U16 - | ptx_parser::ScalarType::E4m3x2 - | ptx_parser::ScalarType::E5m2x2 => unsafe { - LLVMConstInt(llvm_type, u16::from_le_bytes(bytes.try_into()?) as u64, 0) - }, - ptx_parser::ScalarType::S32 - | ptx_parser::ScalarType::B32 - | ptx_parser::ScalarType::U32 => unsafe { - LLVMConstInt(llvm_type, u32::from_le_bytes(bytes.try_into()?) as u64, 0) - }, - ptx_parser::ScalarType::F16 => todo!(), - ptx_parser::ScalarType::BF16 => todo!(), - ptx_parser::ScalarType::U64 => todo!(), - ptx_parser::ScalarType::S64 => todo!(), - ptx_parser::ScalarType::S16x2 => todo!(), - ptx_parser::ScalarType::F32 => todo!(), - ptx_parser::ScalarType::B64 => todo!(), - ptx_parser::ScalarType::F64 => todo!(), - ptx_parser::ScalarType::B128 => todo!(), - ptx_parser::ScalarType::U16x2 => todo!(), - ptx_parser::ScalarType::F16x2 => todo!(), - ptx_parser::ScalarType::BF16x2 => todo!(), - }) + ast::Type::Scalar(scalar) => { + let initializer = match array_init { + [ast::RegOrImmediate::Imm(init)] => init, + _ => return Err(error_mismatched_type()), + }; + get_immediate_value(self.context, scalar, initializer) + } + _ => { + todo!() + } + }; + Ok(initializer) } fn emit_fn_attribute(&self, llvm_object: LLVMValueRef, key: &str, value: &str) { @@ -338,6 +308,20 @@ impl<'a, 'input> ModuleEmitContext<'a, 'input> { } } +fn get_immediate_value( + context: LLVMContextRef, + scalar_type: &ast::ScalarType, + imm: &ast::ImmediateValue, +) -> *mut LLVMValue { + let type_ = get_scalar_type(context, *scalar_type); + match imm { + ast::ImmediateValue::U64(x) => unsafe { LLVMConstInt(type_, *x, 0) }, + ast::ImmediateValue::S64(x) => unsafe { LLVMConstInt(type_, *x as u64, 0) }, + ast::ImmediateValue::F32(x) => unsafe { LLVMConstReal(type_, *x as f64) }, + ast::ImmediateValue::F64(x) => unsafe { LLVMConstReal(type_, *x) }, + } +} + fn llvm_ftz(ftz: bool) -> &'static str { if ftz { "preserve-sign" @@ -404,7 +388,6 @@ impl<'a> MethodEmitContext<'a> { Statement::VectorWrite(vector_write) => self.emit_vector_write(vector_write)?, Statement::SetMode(mode_reg) => self.emit_set_mode(mode_reg)?, Statement::FpSaturate { dst, src, type_ } => self.emit_fp_saturate(type_, dst, src)?, - // No-op Statement::FpModeRequired { .. } => {} }) } @@ -445,7 +428,7 @@ impl<'a> MethodEmitContext<'a> { unsafe { LLVMSetAlignment(alloca, align) }; } if !var.array_init.is_empty() { - todo!() + return Err(error_unreachable()); } Ok(()) } @@ -722,13 +705,7 @@ impl<'a> MethodEmitContext<'a> { } fn emit_constant(&mut self, constant: ConstantDefinition) -> Result<(), TranslateError> { - let type_ = get_scalar_type(self.context, constant.typ); - let value = match constant.value { - ast::ImmediateValue::U64(x) => unsafe { LLVMConstInt(type_, x, 0) }, - ast::ImmediateValue::S64(x) => unsafe { LLVMConstInt(type_, x as u64, 0) }, - ast::ImmediateValue::F32(x) => unsafe { LLVMConstReal(type_, x as f64) }, - ast::ImmediateValue::F64(x) => unsafe { LLVMConstReal(type_, x) }, - }; + let value = get_immediate_value(self.context, &constant.typ, &constant.value); self.resolver.register(constant.dst, value); Ok(()) } diff --git a/ptx/src/pass/mod.rs b/ptx/src/pass/mod.rs index 39bebb0..74bb53c 100644 --- a/ptx/src/pass/mod.rs +++ b/ptx/src/pass/mod.rs @@ -224,7 +224,7 @@ fn error_unknown_symbol>(symbol: T) -> TranslateError { #[cfg(debug_assertions)] fn error_mismatched_type() -> TranslateError { - panic!() + panic!("Mismatched type") } #[cfg(not(debug_assertions))] diff --git a/ptx/src/pass/normalize_identifiers2.rs b/ptx/src/pass/normalize_identifiers2.rs index 05045b7..901f628 100644 --- a/ptx/src/pass/normalize_identifiers2.rs +++ b/ptx/src/pass/normalize_identifiers2.rs @@ -92,7 +92,7 @@ fn run_variable<'input, 'b>( align: variable.align, v_type: variable.v_type, state_space: variable.state_space, - array_init: variable.array_init, + array_init: run_array_init(resolver, &variable.array_init)?, }) } @@ -171,7 +171,7 @@ fn run_multivariable<'input, 'b>( v_type: variable.var.v_type.clone(), state_space: variable.var.state_space, name: ident, - array_init: variable.var.array_init.clone(), + array_init: run_array_init(resolver, &variable.var.array_init)?, })); } } @@ -186,9 +186,22 @@ fn run_multivariable<'input, 'b>( v_type: variable.var.v_type.clone(), state_space: variable.var.state_space, name: ident, - array_init: variable.var.array_init.clone(), + array_init: run_array_init(resolver, &variable.var.array_init)?, })); } } Ok(()) } + +fn run_array_init<'input, 'b>( + resolver: &mut ScopedResolver<'input, 'b>, + array_init: &[ast::RegOrImmediate<&'input str>], +) -> Result>, TranslateError> { + Ok(array_init + .iter() + .map(|elem| match elem { + ast::RegOrImmediate::Reg(name) => Ok(ast::RegOrImmediate::Reg(resolver.get(name)?)), + ast::RegOrImmediate::Imm(imm) => Ok(ast::RegOrImmediate::Imm(*imm)), + }) + .collect::, _>>()?) +} diff --git a/ptx/src/test/ll/const_ident.ll b/ptx/src/test/ll/const_ident.ll new file mode 100644 index 0000000..6d87efd --- /dev/null +++ b/ptx/src/test/ll/const_ident.ll @@ -0,0 +1,55 @@ +@x = addrspace(4) global i64 1 +@y = addrspace(4) global [4 x i64] [i64 4, i64 5, i64 6, i64 0] +@constparams = addrspace(4) global [4 x i64] [i64 ptrtoint (ptr addrspace(4) @x to i64), i64 ptrtoint (ptr addrspace(4) @y to i64), i64 0, i64 0] + +define amdgpu_kernel void @const_ident(ptr addrspace(4) byref(i64) %"49", ptr addrspace(4) byref(i64) %"50") #0 { + %"51" = alloca i64, align 8, addrspace(5) + %"52" = alloca i64, align 8, addrspace(5) + %"53" = alloca i64, align 8, addrspace(5) + %"54" = alloca i64, align 8, addrspace(5) + %"55" = alloca i64, align 8, addrspace(5) + %"56" = alloca i64, align 8, addrspace(5) + %"57" = alloca i64, align 8, addrspace(5) + %"58" = alloca i64, align 8, addrspace(5) + br label %1 + +1: ; preds = %0 + br label %"48" + +"48": ; preds = %1 + %"59" = load i64, ptr addrspace(4) %"49", align 8 + store i64 %"59", ptr addrspace(5) %"51", align 8 + %"60" = load i64, ptr addrspace(4) %"50", align 8 + store i64 %"60", ptr addrspace(5) %"52", align 8 + store i64 ptrtoint (ptr addrspace(4) @x to i64), ptr addrspace(5) %"53", align 8 + store i64 ptrtoint (ptr addrspace(4) @y to i64), ptr addrspace(5) %"54", align 8 + %"63" = load i64, ptr addrspace(4) @constparams, align 8 + store i64 %"63", ptr addrspace(5) %"55", align 8 + %"64" = load i64, ptr addrspace(4) getelementptr inbounds (i8, ptr addrspace(4) @constparams, i64 8), align 8 + store i64 %"64", ptr addrspace(5) %"56", align 8 + %"66" = load i64, ptr addrspace(5) %"53", align 8 + %"67" = load i64, ptr addrspace(5) %"55", align 8 + %"65" = xor i64 %"66", %"67" + store i64 %"65", ptr addrspace(5) %"57", align 8 + %"69" = load i64, ptr addrspace(5) %"54", align 8 + %"70" = load i64, ptr addrspace(5) %"56", align 8 + %"68" = xor i64 %"69", %"70" + store i64 %"68", ptr addrspace(5) %"58", align 8 + %"71" = load i64, ptr addrspace(5) %"52", align 8 + %"72" = load i64, ptr addrspace(5) %"57", align 8 + %"85" = inttoptr i64 %"71" to ptr + store i64 %"72", ptr %"85", align 8 + %"73" = load i64, ptr addrspace(5) %"52", align 8 + %"87" = inttoptr i64 %"73" to ptr + %"45" = getelementptr inbounds i8, ptr %"87", i64 8 + %"74" = load i64, ptr addrspace(5) %"58", align 8 + store i64 %"74", ptr %"45", align 8 + %"75" = load i64, ptr addrspace(5) %"52", align 8 + %"89" = inttoptr i64 %"75" to ptr + %"47" = getelementptr inbounds i8, ptr %"89", i64 8 + %"76" = load i64, ptr addrspace(5) %"58", align 8 + store i64 %"76", ptr %"47", align 8 + ret void +} + +attributes #0 = { "amdgpu-unsafe-fp-atomics"="true" "denormal-fp-math"="preserve-sign" "denormal-fp-math-f32"="preserve-sign" "no-trapping-math"="true" "uniform-work-group-size"="true" } \ No newline at end of file diff --git a/ptx/src/test/ll/global_array_f32.ll b/ptx/src/test/ll/global_array_f32.ll new file mode 100644 index 0000000..201a754 --- /dev/null +++ b/ptx/src/test/ll/global_array_f32.ll @@ -0,0 +1,28 @@ +@foobar = addrspace(1) global [4 x float] [float 1.000000e+00, float 0.000000e+00, float 0.000000e+00, float 0.000000e+00] + +define amdgpu_kernel void @global_array_f32(ptr addrspace(4) byref(i64) %"36", ptr addrspace(4) byref(i64) %"37") #0 { + %"38" = alloca i64, align 8, addrspace(5) + %"39" = alloca i64, align 8, addrspace(5) + %"40" = alloca float, align 4, addrspace(5) + br label %1 + +1: ; preds = %0 + br label %"35" + +"35": ; preds = %1 + store i64 ptrtoint (ptr addrspace(1) @foobar to i64), ptr addrspace(5) %"38", align 8 + %"42" = load i64, ptr addrspace(4) %"37", align 8 + store i64 %"42", ptr addrspace(5) %"39", align 8 + %"43" = load i64, ptr addrspace(5) %"38", align 8 + %"48" = inttoptr i64 %"43" to ptr addrspace(1) + %"34" = getelementptr inbounds i8, ptr addrspace(1) %"48", i64 4 + %"44" = load float, ptr addrspace(1) %"34", align 4 + store float %"44", ptr addrspace(5) %"40", align 4 + %"45" = load i64, ptr addrspace(5) %"39", align 8 + %"46" = load float, ptr addrspace(5) %"40", align 4 + %"49" = inttoptr i64 %"45" to ptr + store float %"46", ptr %"49", align 4 + ret void +} + +attributes #0 = { "amdgpu-unsafe-fp-atomics"="true" "denormal-fp-math"="preserve-sign" "denormal-fp-math-f32"="preserve-sign" "no-trapping-math"="true" "uniform-work-group-size"="true" } diff --git a/ptx/src/test/spirv_run/const_ident.ptx b/ptx/src/test/spirv_run/const_ident.ptx new file mode 100644 index 0000000..9e50e8e --- /dev/null +++ b/ptx/src/test/spirv_run/const_ident.ptx @@ -0,0 +1,39 @@ +.version 6.5 +.target sm_30 +.address_size 64 + +.const .u64 x = 1; +.const .u64 y[4] = {4,5,6}; +.const .u64 constparams[2] = { x, y }; + +.visible .entry const_ident( + .param .u64 input, + .param .u64 output +) +{ + .reg .u64 in_addr; + .reg .u64 out_addr; + .reg .u64 x_addr; + .reg .u64 y_addr; + .reg .u64 constparams_0; + .reg .u64 constparams_1; + .reg .b64 x_equal; + .reg .b64 y_equal; + + ld.param.u64 in_addr, [input]; + ld.param.u64 out_addr, [output]; + + mov.u64 x_addr, x; + mov.u64 y_addr, y; + + ld.const.u64 constparams_0, [constparams+0]; + ld.const.u64 constparams_1, [constparams+8]; + + xor.b64 x_equal, x_addr, constparams_0; + xor.b64 y_equal, y_addr, constparams_1; + + st.u64 [out_addr], x_equal; + st.u64 [out_addr+8], y_equal; + st.u64 [out_addr+8], y_equal; + ret; +} diff --git a/ptx/src/test/spirv_run/global_array_f32.ptx b/ptx/src/test/spirv_run/global_array_f32.ptx new file mode 100644 index 0000000..070fdac --- /dev/null +++ b/ptx/src/test/spirv_run/global_array_f32.ptx @@ -0,0 +1,22 @@ +.version 6.5 +.target sm_30 +.address_size 64 + +.global .f32 foobar[4] = {0f3f800000}; + +.visible .entry global_array_f32( + .param .u64 input, + .param .u64 output +) +{ + .reg .u64 in_addr; + .reg .u64 out_addr; + .reg .f32 temp; + + mov.u64 in_addr, foobar; + ld.param.u64 out_addr, [output]; + + ld.global.f32 temp, [in_addr+4]; + st.f32 [out_addr], temp; + ret; +} \ No newline at end of file diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs index a0758ff..1caf560 100644 --- a/ptx/src/test/spirv_run/mod.rs +++ b/ptx/src/test/spirv_run/mod.rs @@ -144,6 +144,7 @@ test_ptx!(sub, [2u64], [1u64]); test_ptx!(min, [555i32, 444i32], [444i32]); test_ptx!(max, [555i32, 444i32], [555i32]); test_ptx!(global_array, [0xDEADu32], [1u32]); +test_ptx!(global_array_f32, [0x0], [0f32]); test_ptx!(extern_shared, [127u64], [127u64]); test_ptx!(extern_shared_call, [121u64], [123u64]); test_ptx!(rcp, [2f32], [0.5f32]); @@ -261,6 +262,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]); +test_ptx!(const_ident, [0u16], [0u64, 0u64]); test_ptx!(cvt_s16_s8, [0x139231C2u32], [0xFFFFFFC2u32]); test_ptx!(cvt_f64_f32, [0.125f32], [0.125f64]); test_ptx!(prmt, [0x70c507d6u32, 0x6fbd4b5cu32], [0x6fbdd65cu32]); diff --git a/ptx_parser/src/ast.rs b/ptx_parser/src/ast.rs index 3f14540..8285cd3 100644 --- a/ptx_parser/src/ast.rs +++ b/ptx_parser/src/ast.rs @@ -7,6 +7,7 @@ use crate::{ ShuffleMode, VoteMode, }; use bitflags::bitflags; +use derive_more::Display; use std::{alloc::Layout, cmp::Ordering, fmt::Write, num::NonZeroU8}; pub enum Statement { @@ -948,7 +949,7 @@ pub struct Variable { pub v_type: Type, pub state_space: StateSpace, pub name: ID, - pub array_init: Vec, + pub array_init: Vec>, } impl std::fmt::Display for Variable { @@ -1212,6 +1213,12 @@ pub struct ShfDetails { pub mode: FunnelShiftMode, } +#[derive(Clone, Copy, Display)] +pub enum RegOrImmediate { + Reg(Ident), + Imm(ImmediateValue), +} + #[derive(Clone)] pub enum ParsedOperand { Reg(Ident), diff --git a/ptx_parser/src/lib.rs b/ptx_parser/src/lib.rs index 261849e..680a39d 100644 --- a/ptx_parser/src/lib.rs +++ b/ptx_parser/src/lib.rs @@ -332,6 +332,19 @@ fn immediate_value<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult( + stream: &mut PtxParser<'a, 'input>, +) -> PResult> { + trace( + "reg_or_immediate", + alt(( + immediate_value.map(|imm| ast::RegOrImmediate::Imm(imm)), + ident.map(|id| ast::RegOrImmediate::Reg(id)), + )), + ) + .parse_next(stream) +} + pub fn parse_for_errors<'input>(text: &'input str) -> Vec> { let (tokens, mut errors) = lex_with_span_unchecked(text); let parse_result = { @@ -923,9 +936,9 @@ fn multi_variable<'a, 'input: 'a>( let initializer = match state_space { StateSpace::Global | StateSpace::Const => match array_dimensions { Some(ref mut dimensions) => { - opt(array_initializer(vector, type_, dimensions)).parse_next(stream)? + opt(array_initializer(type_, vector, dimensions)).parse_next(stream)? } - None => opt(value_initializer(vector, type_)).parse_next(stream)?, + None => opt(value_initializer(vector)).parse_next(stream)?, }, _ => None, }; @@ -949,10 +962,10 @@ fn multi_variable<'a, 'input: 'a>( } fn array_initializer<'b, 'a: 'b, 'input: 'a>( - vector: Option, type_: ScalarType, + vector: Option, array_dimensions: &'b mut Vec, -) -> impl Parser, Vec, ContextError> + 'b { +) -> impl Parser, Vec>, ContextError> + 'b { trace( "array_initializer", move |stream: &mut PtxParser<'a, 'input>| { @@ -966,15 +979,24 @@ fn array_initializer<'b, 'a: 'b, 'input: 'a>( Token::LBrace, separated::<_, (), (), _, _, _, _>( 0..=array_dimensions[0] as usize, - single_value_append(&mut result, type_), + single_value_append(&mut result), Token::Comma, ), Token::RBrace, ) .parse_next(stream)?; // pad with zeros - let result_size = type_.size_of() as usize * array_dimensions[0] as usize; - result.extend(iter::repeat(0u8).take(result_size - result.len())); + let result_size = array_dimensions[0] as usize; + let default = match type_.kind() { + ScalarKind::Bit | ScalarKind::Unsigned | ScalarKind::Pred => { + ast::ImmediateValue::U64(0) + } + ScalarKind::Signed => ast::ImmediateValue::S64(0), + ScalarKind::Float => ast::ImmediateValue::F64(0.0), + }; + result.extend( + iter::repeat(ast::RegOrImmediate::Imm(default)).take(result_size - result.len()), + ); Ok(result) }, ) @@ -982,8 +1004,7 @@ fn array_initializer<'b, 'a: 'b, 'input: 'a>( fn value_initializer<'a, 'input: 'a>( vector: Option, - type_: ScalarType, -) -> impl Parser, Vec, ContextError> { +) -> impl Parser, Vec>, ContextError> { trace( "value_initializer", move |stream: &mut PtxParser<'a, 'input>| { @@ -993,77 +1014,20 @@ fn value_initializer<'a, 'input: 'a>( if vector.is_some() { return Err(ErrMode::from_error_kind(stream, ErrorKind::Verify)); } - single_value_append(&mut result, type_).parse_next(stream)?; + single_value_append(&mut result).parse_next(stream)?; Ok(result) }, ) } fn single_value_append<'b, 'a: 'b, 'input: 'a>( - accumulator: &'b mut Vec, - type_: ScalarType, + accumulator: &'b mut Vec>, ) -> impl Parser, (), ContextError> + 'b { trace( "single_value_append", move |stream: &mut PtxParser<'a, 'input>| { - let value = immediate_value.parse_next(stream)?; - match (type_, value) { - (ScalarType::U8 | ScalarType::B8, ImmediateValue::U64(x)) => { - accumulator.extend_from_slice(&(x as u8).to_le_bytes()) - } - (ScalarType::U8 | ScalarType::B8, ImmediateValue::S64(x)) => { - accumulator.extend_from_slice(&(x as u8).to_le_bytes()) - } - (ScalarType::U16 | ScalarType::B16, ImmediateValue::U64(x)) => { - accumulator.extend_from_slice(&(x as u16).to_le_bytes()) - } - (ScalarType::U16 | ScalarType::B16, ImmediateValue::S64(x)) => { - accumulator.extend_from_slice(&(x as u16).to_le_bytes()) - } - (ScalarType::U32 | ScalarType::B32, ImmediateValue::U64(x)) => { - accumulator.extend_from_slice(&(x as u32).to_le_bytes()) - } - (ScalarType::U32 | ScalarType::B32, ImmediateValue::S64(x)) => { - accumulator.extend_from_slice(&(x as u32).to_le_bytes()) - } - (ScalarType::U64 | ScalarType::B64, ImmediateValue::U64(x)) => { - accumulator.extend_from_slice(&(x as u64).to_le_bytes()) - } - (ScalarType::U64 | ScalarType::B64, ImmediateValue::S64(x)) => { - accumulator.extend_from_slice(&(x as u64).to_le_bytes()) - } - (ScalarType::S8, ImmediateValue::U64(x)) => { - accumulator.extend_from_slice(&(x as i8).to_le_bytes()) - } - (ScalarType::S8, ImmediateValue::S64(x)) => { - accumulator.extend_from_slice(&(x as i8).to_le_bytes()) - } - (ScalarType::S16, ImmediateValue::U64(x)) => { - accumulator.extend_from_slice(&(x as i16).to_le_bytes()) - } - (ScalarType::S16, ImmediateValue::S64(x)) => { - accumulator.extend_from_slice(&(x as i16).to_le_bytes()) - } - (ScalarType::S32, ImmediateValue::U64(x)) => { - accumulator.extend_from_slice(&(x as i32).to_le_bytes()) - } - (ScalarType::S32, ImmediateValue::S64(x)) => { - accumulator.extend_from_slice(&(x as i32).to_le_bytes()) - } - (ScalarType::S64, ImmediateValue::U64(x)) => { - accumulator.extend_from_slice(&(x as i64).to_le_bytes()) - } - (ScalarType::S64, ImmediateValue::S64(x)) => { - accumulator.extend_from_slice(&(x as i64).to_le_bytes()) - } - (ScalarType::F32, ImmediateValue::F32(x)) => { - accumulator.extend_from_slice(&x.to_le_bytes()) - } - (ScalarType::F64, ImmediateValue::F64(x)) => { - accumulator.extend_from_slice(&x.to_le_bytes()) - } - _ => return Err(ErrMode::from_error_kind(stream, ErrorKind::Verify)), - } + let value = reg_or_immediate.parse_next(stream)?; + accumulator.push(value); Ok(()) }, )