mirror of
https://github.com/vosen/ZLUDA.git
synced 2025-09-10 11:36:42 +00:00
Use Vec<RegOrImmediate> as const/global variable initializer (#490)
This commit is contained in:
Violet
GitHub
parent
b7f3a647d7
commit
e7f10afb51
11 changed files with 255 additions and 148 deletions
|
|
@ -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<Option<ConversionKind>, 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()),
|
||||
}
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -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<u8> 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<SpirvWord>],
|
||||
) -> 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::<Result<Vec<_>, _>>()
|
||||
.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::<Result<Vec<_>, _>>()?;
|
||||
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<LLVMValueRef, TryFromSliceError> {
|
||||
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(())
|
||||
}
|
||||
|
|
|
|||
|
|
@ -224,7 +224,7 @@ fn error_unknown_symbol<T: Into<String>>(symbol: T) -> TranslateError {
|
|||
|
||||
#[cfg(debug_assertions)]
|
||||
fn error_mismatched_type() -> TranslateError {
|
||||
panic!()
|
||||
panic!("Mismatched type")
|
||||
}
|
||||
|
||||
#[cfg(not(debug_assertions))]
|
||||
|
|
|
|||
|
|
@ -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<Vec<ast::RegOrImmediate<SpirvWord>>, 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::<Result<Vec<_>, _>>()?)
|
||||
}
|
||||
|
|
|
|||
55
ptx/src/test/ll/const_ident.ll
Normal file
55
ptx/src/test/ll/const_ident.ll
Normal file
|
|
@ -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" }
|
||||
28
ptx/src/test/ll/global_array_f32.ll
Normal file
28
ptx/src/test/ll/global_array_f32.ll
Normal file
|
|
@ -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" }
|
||||
39
ptx/src/test/spirv_run/const_ident.ptx
Normal file
39
ptx/src/test/spirv_run/const_ident.ptx
Normal file
|
|
@ -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;
|
||||
}
|
||||
22
ptx/src/test/spirv_run/global_array_f32.ptx
Normal file
22
ptx/src/test/spirv_run/global_array_f32.ptx
Normal file
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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]);
|
||||
|
|
|
|||
|
|
@ -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<P: Operand> {
|
||||
|
|
@ -948,7 +949,7 @@ pub struct Variable<ID> {
|
|||
pub v_type: Type,
|
||||
pub state_space: StateSpace,
|
||||
pub name: ID,
|
||||
pub array_init: Vec<u8>,
|
||||
pub array_init: Vec<RegOrImmediate<ID>>,
|
||||
}
|
||||
|
||||
impl<ID: std::fmt::Display> std::fmt::Display for Variable<ID> {
|
||||
|
|
@ -1212,6 +1213,12 @@ pub struct ShfDetails {
|
|||
pub mode: FunnelShiftMode,
|
||||
}
|
||||
|
||||
#[derive(Clone, Copy, Display)]
|
||||
pub enum RegOrImmediate<Ident> {
|
||||
Reg(Ident),
|
||||
Imm(ImmediateValue),
|
||||
}
|
||||
|
||||
#[derive(Clone)]
|
||||
pub enum ParsedOperand<Ident> {
|
||||
Reg(Ident),
|
||||
|
|
|
|||
|
|
@ -332,6 +332,19 @@ fn immediate_value<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<as
|
|||
.parse_next(stream)
|
||||
}
|
||||
|
||||
fn reg_or_immediate<'a, 'input>(
|
||||
stream: &mut PtxParser<'a, 'input>,
|
||||
) -> PResult<ast::RegOrImmediate<&'input str>> {
|
||||
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<PtxError<'input>> {
|
||||
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<NonZeroU8>,
|
||||
type_: ScalarType,
|
||||
vector: Option<NonZeroU8>,
|
||||
array_dimensions: &'b mut Vec<u32>,
|
||||
) -> impl Parser<PtxParser<'a, 'input>, Vec<u8>, ContextError> + 'b {
|
||||
) -> impl Parser<PtxParser<'a, 'input>, Vec<RegOrImmediate<&'input str>>, 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<NonZeroU8>,
|
||||
type_: ScalarType,
|
||||
) -> impl Parser<PtxParser<'a, 'input>, Vec<u8>, ContextError> {
|
||||
) -> impl Parser<PtxParser<'a, 'input>, Vec<RegOrImmediate<&'input str>>, 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<u8>,
|
||||
type_: ScalarType,
|
||||
accumulator: &'b mut Vec<RegOrImmediate<&'input str>>,
|
||||
) -> impl Parser<PtxParser<'a, 'input>, (), 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(())
|
||||
},
|
||||
)
|
||||
|
|
|
|||
Loading…
Add table
Add a link
Reference in a new issue