Support float constants
This commit is contained in:
parent
17b788f2a7
commit
40bdb83e6b
|
@ -621,17 +621,25 @@ pub struct Arg5<P: ArgParams> {
|
||||||
pub src3: P::Operand,
|
pub src3: P::Operand,
|
||||||
}
|
}
|
||||||
|
|
||||||
|
#[derive(Copy, Clone)]
|
||||||
|
pub enum ImmediateValue {
|
||||||
|
U64(u64),
|
||||||
|
S64(i64),
|
||||||
|
F32(f32),
|
||||||
|
F64(f64),
|
||||||
|
}
|
||||||
|
|
||||||
#[derive(Copy, Clone)]
|
#[derive(Copy, Clone)]
|
||||||
pub enum Operand<ID> {
|
pub enum Operand<ID> {
|
||||||
Reg(ID),
|
Reg(ID),
|
||||||
RegOffset(ID, i32),
|
RegOffset(ID, i32),
|
||||||
Imm(u32),
|
Imm(ImmediateValue),
|
||||||
}
|
}
|
||||||
|
|
||||||
#[derive(Copy, Clone)]
|
#[derive(Copy, Clone)]
|
||||||
pub enum CallOperand<ID> {
|
pub enum CallOperand<ID> {
|
||||||
Reg(ID),
|
Reg(ID),
|
||||||
Imm(u32),
|
Imm(ImmediateValue),
|
||||||
}
|
}
|
||||||
|
|
||||||
pub enum IdOrVector<ID> {
|
pub enum IdOrVector<ID> {
|
||||||
|
@ -642,7 +650,7 @@ pub enum IdOrVector<ID> {
|
||||||
pub enum OperandOrVector<ID> {
|
pub enum OperandOrVector<ID> {
|
||||||
Reg(ID),
|
Reg(ID),
|
||||||
RegOffset(ID, i32),
|
RegOffset(ID, i32),
|
||||||
Imm(u32),
|
Imm(ImmediateValue),
|
||||||
Vec(Vec<ID>),
|
Vec(Vec<ID>),
|
||||||
}
|
}
|
||||||
|
|
||||||
|
@ -1028,7 +1036,7 @@ pub struct MinMaxFloat {
|
||||||
}
|
}
|
||||||
|
|
||||||
pub enum NumsOrArrays<'a> {
|
pub enum NumsOrArrays<'a> {
|
||||||
Nums(Vec<&'a str>),
|
Nums(Vec<(&'a str, u32)>),
|
||||||
Arrays(Vec<NumsOrArrays<'a>>),
|
Arrays(Vec<NumsOrArrays<'a>>),
|
||||||
}
|
}
|
||||||
|
|
||||||
|
@ -1076,8 +1084,8 @@ impl<'a> NumsOrArrays<'a> {
|
||||||
if vec.len() > *dim as usize {
|
if vec.len() > *dim as usize {
|
||||||
return Err(PtxError::ZeroDimensionArray);
|
return Err(PtxError::ZeroDimensionArray);
|
||||||
}
|
}
|
||||||
for (idx, val) in vec.iter().enumerate() {
|
for (idx, (val, radix)) in vec.iter().enumerate() {
|
||||||
Self::parse_and_copy_single(t, idx, val, result)?;
|
Self::parse_and_copy_single(t, idx, val, *radix, result)?;
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
NumsOrArrays::Arrays(_) => return Err(PtxError::ZeroDimensionArray),
|
NumsOrArrays::Arrays(_) => return Err(PtxError::ZeroDimensionArray),
|
||||||
|
@ -1107,42 +1115,43 @@ impl<'a> NumsOrArrays<'a> {
|
||||||
t: SizedScalarType,
|
t: SizedScalarType,
|
||||||
idx: usize,
|
idx: usize,
|
||||||
str_val: &str,
|
str_val: &str,
|
||||||
|
radix: u32,
|
||||||
output: &mut [u8],
|
output: &mut [u8],
|
||||||
) -> Result<(), PtxError> {
|
) -> Result<(), PtxError> {
|
||||||
match t {
|
match t {
|
||||||
SizedScalarType::B8 | SizedScalarType::U8 => {
|
SizedScalarType::B8 | SizedScalarType::U8 => {
|
||||||
Self::parse_and_copy_single_t::<u8>(idx, str_val, output)?;
|
Self::parse_and_copy_single_t::<u8>(idx, str_val, radix, output)?;
|
||||||
}
|
}
|
||||||
SizedScalarType::B16 | SizedScalarType::U16 => {
|
SizedScalarType::B16 | SizedScalarType::U16 => {
|
||||||
Self::parse_and_copy_single_t::<u16>(idx, str_val, output)?;
|
Self::parse_and_copy_single_t::<u16>(idx, str_val, radix, output)?;
|
||||||
}
|
}
|
||||||
SizedScalarType::B32 | SizedScalarType::U32 => {
|
SizedScalarType::B32 | SizedScalarType::U32 => {
|
||||||
Self::parse_and_copy_single_t::<u32>(idx, str_val, output)?;
|
Self::parse_and_copy_single_t::<u32>(idx, str_val, radix, output)?;
|
||||||
}
|
}
|
||||||
SizedScalarType::B64 | SizedScalarType::U64 => {
|
SizedScalarType::B64 | SizedScalarType::U64 => {
|
||||||
Self::parse_and_copy_single_t::<u64>(idx, str_val, output)?;
|
Self::parse_and_copy_single_t::<u64>(idx, str_val, radix, output)?;
|
||||||
}
|
}
|
||||||
SizedScalarType::S8 => {
|
SizedScalarType::S8 => {
|
||||||
Self::parse_and_copy_single_t::<i8>(idx, str_val, output)?;
|
Self::parse_and_copy_single_t::<i8>(idx, str_val, radix, output)?;
|
||||||
}
|
}
|
||||||
SizedScalarType::S16 => {
|
SizedScalarType::S16 => {
|
||||||
Self::parse_and_copy_single_t::<i16>(idx, str_val, output)?;
|
Self::parse_and_copy_single_t::<i16>(idx, str_val, radix, output)?;
|
||||||
}
|
}
|
||||||
SizedScalarType::S32 => {
|
SizedScalarType::S32 => {
|
||||||
Self::parse_and_copy_single_t::<i32>(idx, str_val, output)?;
|
Self::parse_and_copy_single_t::<i32>(idx, str_val, radix, output)?;
|
||||||
}
|
}
|
||||||
SizedScalarType::S64 => {
|
SizedScalarType::S64 => {
|
||||||
Self::parse_and_copy_single_t::<i64>(idx, str_val, output)?;
|
Self::parse_and_copy_single_t::<i64>(idx, str_val, radix, output)?;
|
||||||
}
|
}
|
||||||
SizedScalarType::F16 => {
|
SizedScalarType::F16 => {
|
||||||
Self::parse_and_copy_single_t::<f16>(idx, str_val, output)?;
|
Self::parse_and_copy_single_t::<f16>(idx, str_val, radix, output)?;
|
||||||
}
|
}
|
||||||
SizedScalarType::F16x2 => todo!(),
|
SizedScalarType::F16x2 => todo!(),
|
||||||
SizedScalarType::F32 => {
|
SizedScalarType::F32 => {
|
||||||
Self::parse_and_copy_single_t::<f32>(idx, str_val, output)?;
|
Self::parse_and_copy_single_t::<f32>(idx, str_val, radix, output)?;
|
||||||
}
|
}
|
||||||
SizedScalarType::F64 => {
|
SizedScalarType::F64 => {
|
||||||
Self::parse_and_copy_single_t::<f64>(idx, str_val, output)?;
|
Self::parse_and_copy_single_t::<f64>(idx, str_val, radix, output)?;
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
Ok(())
|
Ok(())
|
||||||
|
@ -1151,6 +1160,7 @@ impl<'a> NumsOrArrays<'a> {
|
||||||
fn parse_and_copy_single_t<T: Copy + FromStr>(
|
fn parse_and_copy_single_t<T: Copy + FromStr>(
|
||||||
idx: usize,
|
idx: usize,
|
||||||
str_val: &str,
|
str_val: &str,
|
||||||
|
_radix: u32, // TODO: use this to properly support hex literals
|
||||||
output: &mut [u8],
|
output: &mut [u8],
|
||||||
) -> Result<(), PtxError>
|
) -> Result<(), PtxError>
|
||||||
where
|
where
|
||||||
|
@ -1200,8 +1210,8 @@ mod tests {
|
||||||
#[test]
|
#[test]
|
||||||
fn array_auto_sizes_0_dimension() {
|
fn array_auto_sizes_0_dimension() {
|
||||||
let inp = NumsOrArrays::Arrays(vec![
|
let inp = NumsOrArrays::Arrays(vec![
|
||||||
NumsOrArrays::Nums(vec!["1", "2"]),
|
NumsOrArrays::Nums(vec![("1", 10), ("2", 10)]),
|
||||||
NumsOrArrays::Nums(vec!["3", "4"]),
|
NumsOrArrays::Nums(vec![("3", 10), ("4", 10)]),
|
||||||
]);
|
]);
|
||||||
let mut dimensions = vec![0u32, 2];
|
let mut dimensions = vec![0u32, 2];
|
||||||
assert_eq!(
|
assert_eq!(
|
||||||
|
@ -1214,8 +1224,8 @@ mod tests {
|
||||||
#[test]
|
#[test]
|
||||||
fn array_fails_wrong_structure() {
|
fn array_fails_wrong_structure() {
|
||||||
let inp = NumsOrArrays::Arrays(vec![
|
let inp = NumsOrArrays::Arrays(vec![
|
||||||
NumsOrArrays::Nums(vec!["1", "2"]),
|
NumsOrArrays::Nums(vec![("1", 10), ("2", 10)]),
|
||||||
NumsOrArrays::Arrays(vec![NumsOrArrays::Nums(vec!["1"])]),
|
NumsOrArrays::Arrays(vec![NumsOrArrays::Nums(vec![("1", 10)])]),
|
||||||
]);
|
]);
|
||||||
let mut dimensions = vec![0u32, 2];
|
let mut dimensions = vec![0u32, 2];
|
||||||
assert!(inp.to_vec(SizedScalarType::B8, &mut dimensions).is_err());
|
assert!(inp.to_vec(SizedScalarType::B8, &mut dimensions).is_err());
|
||||||
|
@ -1224,8 +1234,8 @@ mod tests {
|
||||||
#[test]
|
#[test]
|
||||||
fn array_fails_too_long_component() {
|
fn array_fails_too_long_component() {
|
||||||
let inp = NumsOrArrays::Arrays(vec![
|
let inp = NumsOrArrays::Arrays(vec![
|
||||||
NumsOrArrays::Nums(vec!["1", "2", "3"]),
|
NumsOrArrays::Nums(vec![("1", 10), ("2", 10), ("3", 10)]),
|
||||||
NumsOrArrays::Nums(vec!["4", "5"]),
|
NumsOrArrays::Nums(vec![("4", 10), ("5", 10)]),
|
||||||
]);
|
]);
|
||||||
let mut dimensions = vec![0u32, 2];
|
let mut dimensions = vec![0u32, 2];
|
||||||
assert!(inp.to_vec(SizedScalarType::B8, &mut dimensions).is_err());
|
assert!(inp.to_vec(SizedScalarType::B8, &mut dimensions).is_err());
|
||||||
|
|
|
@ -15,12 +15,16 @@ match {
|
||||||
r"\s+" => { },
|
r"\s+" => { },
|
||||||
r"//[^\n\r]*[\n\r]*" => { },
|
r"//[^\n\r]*[\n\r]*" => { },
|
||||||
r"/\*([^\*]*\*+[^\*/])*([^\*]*\*+|[^\*])*\*/" => { },
|
r"/\*([^\*]*\*+[^\*/])*([^\*]*\*+|[^\*])*\*/" => { },
|
||||||
r"-?[?:0x]?[0-9]+" => Num,
|
r"0[fF][0-9a-zA-Z]{8}" => F32NumToken,
|
||||||
|
r"0[dD][0-9a-zA-Z]{16}" => F64NumToken,
|
||||||
|
r"0[xX][0-9a-zA-Z]+U?" => HexNumToken,
|
||||||
|
r"[0-9]+U?" => DecimalNumToken,
|
||||||
r#""[^"]*""# => String,
|
r#""[^"]*""# => String,
|
||||||
r"[0-9]+\.[0-9]+" => VersionNumber,
|
r"[0-9]+\.[0-9]+" => VersionNumber,
|
||||||
"!",
|
"!",
|
||||||
"(", ")",
|
"(", ")",
|
||||||
"+",
|
"+",
|
||||||
|
"-",
|
||||||
",",
|
",",
|
||||||
".",
|
".",
|
||||||
":",
|
":",
|
||||||
|
@ -181,6 +185,74 @@ ExtendedID : &'input str = {
|
||||||
ID
|
ID
|
||||||
}
|
}
|
||||||
|
|
||||||
|
NumToken: (&'input str, u32, bool) = {
|
||||||
|
<s:HexNumToken> => {
|
||||||
|
if s.ends_with('U') {
|
||||||
|
(&s[2..s.len() - 1], 16, true)
|
||||||
|
} else {
|
||||||
|
(&s[2..], 16, false)
|
||||||
|
}
|
||||||
|
},
|
||||||
|
<s:DecimalNumToken> => {
|
||||||
|
let radix = if s.starts_with('0') { 8 } else { 10 };
|
||||||
|
if s.ends_with('U') {
|
||||||
|
(&s[..s.len() - 1], radix, true)
|
||||||
|
} else {
|
||||||
|
(s, radix, false)
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
F32Num: f32 = {
|
||||||
|
<s:F32NumToken> =>? {
|
||||||
|
match u32::from_str_radix(&s[2..], 16) {
|
||||||
|
Ok(x) => Ok(unsafe { std::mem::transmute::<_, f32>(x) }),
|
||||||
|
Err(err) => Err(ParseError::User { error: ast::PtxError::from(err) })
|
||||||
|
}
|
||||||
|
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
F64Num: f64 = {
|
||||||
|
<s:F64NumToken> =>? {
|
||||||
|
match u64::from_str_radix(&s[2..], 16) {
|
||||||
|
Ok(x) => Ok(unsafe { std::mem::transmute::<_, f64>(x) }),
|
||||||
|
Err(err) => Err(ParseError::User { error: ast::PtxError::from(err) })
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
U8Num: u8 = {
|
||||||
|
<x:NumToken> =>? {
|
||||||
|
let (text, radix, _) = x;
|
||||||
|
match u8::from_str_radix(text, radix) {
|
||||||
|
Ok(x) => Ok(x),
|
||||||
|
Err(err) => Err(ParseError::User { error: ast::PtxError::from(err) })
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
U32Num: u32 = {
|
||||||
|
<x:NumToken> =>? {
|
||||||
|
let (text, radix, _) = x;
|
||||||
|
match u32::from_str_radix(text, radix) {
|
||||||
|
Ok(x) => Ok(x),
|
||||||
|
Err(err) => Err(ParseError::User { error: ast::PtxError::from(err) })
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
// TODO: handle negative number properly
|
||||||
|
S32Num: i32 = {
|
||||||
|
<sign:"-"?> <x:NumToken> =>? {
|
||||||
|
let (text, radix, _) = x;
|
||||||
|
match i32::from_str_radix(text, radix) {
|
||||||
|
Ok(x) => Ok(if sign.is_some() { -x } else { x }),
|
||||||
|
Err(err) => Err(ParseError::User { error: ast::PtxError::from(err) })
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
pub Module: ast::Module<'input> = {
|
pub Module: ast::Module<'input> = {
|
||||||
<v:Version> Target <d:Directive*> => {
|
<v:Version> Target <d:Directive*> => {
|
||||||
ast::Module { version: v, directives: without_none(d) }
|
ast::Module { version: v, directives: without_none(d) }
|
||||||
|
@ -218,7 +290,7 @@ Directive: Option<ast::Directive<'input, ast::ParsedArgParams<'input>>> = {
|
||||||
};
|
};
|
||||||
|
|
||||||
AddressSize = {
|
AddressSize = {
|
||||||
".address_size" Num
|
".address_size" U8Num
|
||||||
};
|
};
|
||||||
|
|
||||||
Function: ast::Function<'input, &'input str, ast::Statement<ast::ParsedArgParams<'input>>> = {
|
Function: ast::Function<'input, &'input str, ast::Statement<ast::ParsedArgParams<'input>>> = {
|
||||||
|
@ -328,7 +400,7 @@ DebugDirective: () = {
|
||||||
|
|
||||||
// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#debugging-directives-loc
|
// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#debugging-directives-loc
|
||||||
DebugLocation = {
|
DebugLocation = {
|
||||||
".loc" Num Num Num
|
".loc" U32Num U32Num U32Num
|
||||||
};
|
};
|
||||||
|
|
||||||
Label: &'input str = {
|
Label: &'input str = {
|
||||||
|
@ -336,10 +408,7 @@ Label: &'input str = {
|
||||||
};
|
};
|
||||||
|
|
||||||
Align: u32 = {
|
Align: u32 = {
|
||||||
".align" <a:Num> => {
|
".align" <x:U32Num> => x
|
||||||
let align = a.parse::<u32>();
|
|
||||||
align.unwrap_with(errors)
|
|
||||||
}
|
|
||||||
};
|
};
|
||||||
|
|
||||||
// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parameterized-variable-names
|
// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parameterized-variable-names
|
||||||
|
@ -348,10 +417,7 @@ MultiVariable: ast::MultiVariable<&'input str> = {
|
||||||
}
|
}
|
||||||
|
|
||||||
VariableParam: u32 = {
|
VariableParam: u32 = {
|
||||||
"<" <n:Num> ">" => {
|
"<" <n:U32Num> ">" => n
|
||||||
let size = n.parse::<u32>();
|
|
||||||
size.unwrap_with(errors)
|
|
||||||
}
|
|
||||||
}
|
}
|
||||||
|
|
||||||
Variable: ast::Variable<ast::VariableType, &'input str> = {
|
Variable: ast::Variable<ast::VariableType, &'input str> = {
|
||||||
|
@ -1239,29 +1305,51 @@ ArithFloat: ast::ArithFloat = {
|
||||||
|
|
||||||
Operand: ast::Operand<&'input str> = {
|
Operand: ast::Operand<&'input str> = {
|
||||||
<r:ExtendedID> => ast::Operand::Reg(r),
|
<r:ExtendedID> => ast::Operand::Reg(r),
|
||||||
<r:ExtendedID> "+" <o:Num> => {
|
<r:ExtendedID> "+" <offset:S32Num> => ast::Operand::RegOffset(r, offset),
|
||||||
let offset = o.parse::<i32>();
|
<x:ImmediateValue> => ast::Operand::Imm(x)
|
||||||
let offset = offset.unwrap_with(errors);
|
|
||||||
ast::Operand::RegOffset(r, offset)
|
|
||||||
},
|
|
||||||
// TODO: start parsing whole constants sub-language:
|
|
||||||
// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#constants
|
|
||||||
<o:Num> => {
|
|
||||||
let offset = o.parse::<u32>();
|
|
||||||
let offset = offset.unwrap_with(errors);
|
|
||||||
ast::Operand::Imm(offset)
|
|
||||||
}
|
|
||||||
};
|
};
|
||||||
|
|
||||||
CallOperand: ast::CallOperand<&'input str> = {
|
CallOperand: ast::CallOperand<&'input str> = {
|
||||||
<r:ExtendedID> => ast::CallOperand::Reg(r),
|
<r:ExtendedID> => ast::CallOperand::Reg(r),
|
||||||
<o:Num> => {
|
<x:ImmediateValue> => ast::CallOperand::Imm(x)
|
||||||
let offset = o.parse::<u32>();
|
|
||||||
let offset = offset.unwrap_with(errors);
|
|
||||||
ast::CallOperand::Imm(offset)
|
|
||||||
}
|
|
||||||
};
|
};
|
||||||
|
|
||||||
|
// TODO: start parsing whole constants sub-language:
|
||||||
|
// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#constants
|
||||||
|
ImmediateValue: ast::ImmediateValue = {
|
||||||
|
// TODO: treat negation correctly
|
||||||
|
<neg:"-"?> <x:NumToken> =>? {
|
||||||
|
let (num, radix, is_unsigned) = x;
|
||||||
|
if neg.is_some() {
|
||||||
|
match i64::from_str_radix(num, radix) {
|
||||||
|
Ok(x) => Ok(ast::ImmediateValue::S64(-x)),
|
||||||
|
Err(err) => Err(ParseError::User { error: ast::PtxError::ParseInt(err) })
|
||||||
|
}
|
||||||
|
} else if is_unsigned {
|
||||||
|
match u64::from_str_radix(num, radix) {
|
||||||
|
Ok(x) => Ok(ast::ImmediateValue::U64(x)),
|
||||||
|
Err(err) => Err(ParseError::User { error: ast::PtxError::ParseInt(err) })
|
||||||
|
}
|
||||||
|
} else {
|
||||||
|
match i64::from_str_radix(num, radix) {
|
||||||
|
Ok(x) => Ok(ast::ImmediateValue::S64(x)),
|
||||||
|
Err(_) => {
|
||||||
|
match u64::from_str_radix(num, radix) {
|
||||||
|
Ok(x) => Ok(ast::ImmediateValue::U64(x)),
|
||||||
|
Err(err) => Err(ParseError::User { error: ast::PtxError::ParseInt(err) })
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
},
|
||||||
|
<f:F32Num> => {
|
||||||
|
ast::ImmediateValue::F32(f)
|
||||||
|
},
|
||||||
|
<f:F64Num> => {
|
||||||
|
ast::ImmediateValue::F64(f)
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
Arg1: ast::Arg1<ast::ParsedArgParams<'input>> = {
|
Arg1: ast::Arg1<ast::ParsedArgParams<'input>> = {
|
||||||
<src:ExtendedID> => ast::Arg1{<>}
|
<src:ExtendedID> => ast::Arg1{<>}
|
||||||
};
|
};
|
||||||
|
@ -1332,7 +1420,7 @@ VectorPrefix: u8 = {
|
||||||
|
|
||||||
// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#debugging-directives-file
|
// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#debugging-directives-file
|
||||||
File = {
|
File = {
|
||||||
".file" Num String ("," Num "," Num)?
|
".file" U32Num String ("," U32Num "," U32Num)?
|
||||||
};
|
};
|
||||||
|
|
||||||
// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#debugging-directives-section
|
// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#debugging-directives-section
|
||||||
|
@ -1341,11 +1429,11 @@ Section = {
|
||||||
};
|
};
|
||||||
|
|
||||||
SectionDwarfLines: () = {
|
SectionDwarfLines: () = {
|
||||||
BitType Comma<Num>,
|
BitType Comma<U32Num>,
|
||||||
".b32" SectionLabel,
|
".b32" SectionLabel,
|
||||||
".b64" SectionLabel,
|
".b64" SectionLabel,
|
||||||
".b32" SectionLabel "+" Num,
|
".b32" SectionLabel "+" U32Num,
|
||||||
".b64" SectionLabel "+" Num,
|
".b64" SectionLabel "+" U32Num,
|
||||||
};
|
};
|
||||||
|
|
||||||
SectionLabel = {
|
SectionLabel = {
|
||||||
|
@ -1409,9 +1497,7 @@ ArrayEmptyDimension = {
|
||||||
}
|
}
|
||||||
|
|
||||||
ArrayDimension: u32 = {
|
ArrayDimension: u32 = {
|
||||||
"[" <n:Num> "]" =>? {
|
"[" <n:U32Num> "]" => n,
|
||||||
str::parse::<u32>(n).map_err(|e| ParseError::User { error: ast::PtxError::ParseInt(e) })
|
|
||||||
}
|
|
||||||
}
|
}
|
||||||
|
|
||||||
ArrayInitializer: ast::NumsOrArrays<'input> = {
|
ArrayInitializer: ast::NumsOrArrays<'input> = {
|
||||||
|
@ -1424,7 +1510,7 @@ NumsOrArraysBracket: ast::NumsOrArrays<'input> = {
|
||||||
|
|
||||||
NumsOrArrays: ast::NumsOrArrays<'input> = {
|
NumsOrArrays: ast::NumsOrArrays<'input> = {
|
||||||
<n:Comma<NumsOrArraysBracket>> => ast::NumsOrArrays::Arrays(n),
|
<n:Comma<NumsOrArraysBracket>> => ast::NumsOrArrays::Arrays(n),
|
||||||
<n:CommaNonEmpty<Num>> => ast::NumsOrArrays::Nums(n),
|
<n:CommaNonEmpty<NumToken>> => ast::NumsOrArrays::Nums(n.into_iter().map(|(x,radix,_)| (x, radix)).collect()),
|
||||||
}
|
}
|
||||||
|
|
||||||
Comma<T>: Vec<T> = {
|
Comma<T>: Vec<T> = {
|
||||||
|
|
21
ptx/src/test/spirv_run/constant_f32.ptx
Normal file
21
ptx/src/test/spirv_run/constant_f32.ptx
Normal file
|
@ -0,0 +1,21 @@
|
||||||
|
.version 6.5
|
||||||
|
.target sm_30
|
||||||
|
.address_size 64
|
||||||
|
|
||||||
|
.visible .entry constant_f32(
|
||||||
|
.param .u64 input,
|
||||||
|
.param .u64 output
|
||||||
|
)
|
||||||
|
{
|
||||||
|
.reg .u64 in_addr;
|
||||||
|
.reg .u64 out_addr;
|
||||||
|
.reg .f32 temp;
|
||||||
|
|
||||||
|
ld.param.u64 in_addr, [input];
|
||||||
|
ld.param.u64 out_addr, [output];
|
||||||
|
|
||||||
|
ld.f32 temp, [in_addr];
|
||||||
|
mul.f32 temp, temp, 0f3f000000; // 0.5
|
||||||
|
st.f32 [out_addr], temp;
|
||||||
|
ret;
|
||||||
|
}
|
57
ptx/src/test/spirv_run/constant_f32.spvtxt
Normal file
57
ptx/src/test/spirv_run/constant_f32.spvtxt
Normal file
|
@ -0,0 +1,57 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.3
|
||||||
|
; Generator: rspirv
|
||||||
|
; Bound: 32
|
||||||
|
OpCapability GenericPointer
|
||||||
|
OpCapability Linkage
|
||||||
|
OpCapability Addresses
|
||||||
|
OpCapability Kernel
|
||||||
|
OpCapability Int8
|
||||||
|
OpCapability Int16
|
||||||
|
OpCapability Int64
|
||||||
|
OpCapability Float16
|
||||||
|
OpCapability Float64
|
||||||
|
OpCapability FunctionFloatControlINTEL
|
||||||
|
OpExtension "SPV_INTEL_float_controls2"
|
||||||
|
%24 = OpExtInstImport "OpenCL.std"
|
||||||
|
OpMemoryModel Physical64 OpenCL
|
||||||
|
OpEntryPoint Kernel %1 "constant_f32"
|
||||||
|
OpDecorate %1 FunctionDenormModeINTEL 32 Preserve
|
||||||
|
%25 = OpTypeVoid
|
||||||
|
%26 = OpTypeInt 64 0
|
||||||
|
%27 = OpTypeFunction %25 %26 %26
|
||||||
|
%28 = OpTypePointer Function %26
|
||||||
|
%29 = OpTypeFloat 32
|
||||||
|
%30 = OpTypePointer Function %29
|
||||||
|
%31 = OpTypePointer Generic %29
|
||||||
|
%19 = OpConstant %29 0.5
|
||||||
|
%1 = OpFunction %25 None %27
|
||||||
|
%7 = OpFunctionParameter %26
|
||||||
|
%8 = OpFunctionParameter %26
|
||||||
|
%22 = OpLabel
|
||||||
|
%2 = OpVariable %28 Function
|
||||||
|
%3 = OpVariable %28 Function
|
||||||
|
%4 = OpVariable %28 Function
|
||||||
|
%5 = OpVariable %28 Function
|
||||||
|
%6 = OpVariable %30 Function
|
||||||
|
OpStore %2 %7
|
||||||
|
OpStore %3 %8
|
||||||
|
%10 = OpLoad %26 %2
|
||||||
|
%9 = OpCopyObject %26 %10
|
||||||
|
OpStore %4 %9
|
||||||
|
%12 = OpLoad %26 %3
|
||||||
|
%11 = OpCopyObject %26 %12
|
||||||
|
OpStore %5 %11
|
||||||
|
%14 = OpLoad %26 %4
|
||||||
|
%20 = OpConvertUToPtr %31 %14
|
||||||
|
%13 = OpLoad %29 %20
|
||||||
|
OpStore %6 %13
|
||||||
|
%16 = OpLoad %29 %6
|
||||||
|
%15 = OpFMul %29 %16 %19
|
||||||
|
OpStore %6 %15
|
||||||
|
%17 = OpLoad %26 %5
|
||||||
|
%18 = OpLoad %29 %6
|
||||||
|
%21 = OpConvertUToPtr %31 %17
|
||||||
|
OpStore %21 %18
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
21
ptx/src/test/spirv_run/constant_negative.ptx
Normal file
21
ptx/src/test/spirv_run/constant_negative.ptx
Normal file
|
@ -0,0 +1,21 @@
|
||||||
|
.version 6.5
|
||||||
|
.target sm_30
|
||||||
|
.address_size 64
|
||||||
|
|
||||||
|
.visible .entry constant_negative(
|
||||||
|
.param .u64 input,
|
||||||
|
.param .u64 output
|
||||||
|
)
|
||||||
|
{
|
||||||
|
.reg .u64 in_addr;
|
||||||
|
.reg .u64 out_addr;
|
||||||
|
.reg .s32 temp;
|
||||||
|
|
||||||
|
ld.param.u64 in_addr, [input];
|
||||||
|
ld.param.u64 out_addr, [output];
|
||||||
|
|
||||||
|
ld.s32 temp, [in_addr];
|
||||||
|
mul.lo.s32 temp, temp, -1;
|
||||||
|
st.s32 [out_addr], temp;
|
||||||
|
ret;
|
||||||
|
}
|
56
ptx/src/test/spirv_run/constant_negative.spvtxt
Normal file
56
ptx/src/test/spirv_run/constant_negative.spvtxt
Normal file
|
@ -0,0 +1,56 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.3
|
||||||
|
; Generator: rspirv
|
||||||
|
; Bound: 32
|
||||||
|
OpCapability GenericPointer
|
||||||
|
OpCapability Linkage
|
||||||
|
OpCapability Addresses
|
||||||
|
OpCapability Kernel
|
||||||
|
OpCapability Int8
|
||||||
|
OpCapability Int16
|
||||||
|
OpCapability Int64
|
||||||
|
OpCapability Float16
|
||||||
|
OpCapability Float64
|
||||||
|
OpCapability FunctionFloatControlINTEL
|
||||||
|
OpExtension "SPV_INTEL_float_controls2"
|
||||||
|
%24 = OpExtInstImport "OpenCL.std"
|
||||||
|
OpMemoryModel Physical64 OpenCL
|
||||||
|
OpEntryPoint Kernel %1 "constant_negative"
|
||||||
|
%25 = OpTypeVoid
|
||||||
|
%26 = OpTypeInt 64 0
|
||||||
|
%27 = OpTypeFunction %25 %26 %26
|
||||||
|
%28 = OpTypePointer Function %26
|
||||||
|
%29 = OpTypeInt 32 0
|
||||||
|
%30 = OpTypePointer Function %29
|
||||||
|
%31 = OpTypePointer Generic %29
|
||||||
|
%19 = OpConstant %29 4294967295
|
||||||
|
%1 = OpFunction %25 None %27
|
||||||
|
%7 = OpFunctionParameter %26
|
||||||
|
%8 = OpFunctionParameter %26
|
||||||
|
%22 = OpLabel
|
||||||
|
%2 = OpVariable %28 Function
|
||||||
|
%3 = OpVariable %28 Function
|
||||||
|
%4 = OpVariable %28 Function
|
||||||
|
%5 = OpVariable %28 Function
|
||||||
|
%6 = OpVariable %30 Function
|
||||||
|
OpStore %2 %7
|
||||||
|
OpStore %3 %8
|
||||||
|
%10 = OpLoad %26 %2
|
||||||
|
%9 = OpCopyObject %26 %10
|
||||||
|
OpStore %4 %9
|
||||||
|
%12 = OpLoad %26 %3
|
||||||
|
%11 = OpCopyObject %26 %12
|
||||||
|
OpStore %5 %11
|
||||||
|
%14 = OpLoad %26 %4
|
||||||
|
%20 = OpConvertUToPtr %31 %14
|
||||||
|
%13 = OpLoad %29 %20
|
||||||
|
OpStore %6 %13
|
||||||
|
%16 = OpLoad %29 %6
|
||||||
|
%15 = OpIMul %29 %16 %19
|
||||||
|
OpStore %6 %15
|
||||||
|
%17 = OpLoad %26 %5
|
||||||
|
%18 = OpLoad %29 %6
|
||||||
|
%21 = OpConvertUToPtr %31 %17
|
||||||
|
OpStore %21 %18
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
|
@ -87,6 +87,8 @@ test_ptx!(rcp, [2f32], [0.5f32]);
|
||||||
// TODO: mul_ftz fails because IGC does not yet handle SPV_INTEL_float_controls2
|
// TODO: mul_ftz fails because IGC does not yet handle SPV_INTEL_float_controls2
|
||||||
// test_ptx!(mul_ftz, [0b1_00000000_10000000000000000000000u32, 0x3f000000u32], [0u32]);
|
// test_ptx!(mul_ftz, [0b1_00000000_10000000000000000000000u32, 0x3f000000u32], [0u32]);
|
||||||
test_ptx!(mul_non_ftz, [0b1_00000000_10000000000000000000000u32, 0x3f000000u32], [0b1_00000000_01000000000000000000000u32]);
|
test_ptx!(mul_non_ftz, [0b1_00000000_10000000000000000000000u32, 0x3f000000u32], [0b1_00000000_01000000000000000000000u32]);
|
||||||
|
test_ptx!(constant_f32, [10f32], [5f32]);
|
||||||
|
test_ptx!(constant_negative, [-101i32], [101i32]);
|
||||||
|
|
||||||
struct DisplayError<T: Debug> {
|
struct DisplayError<T: Debug> {
|
||||||
err: T,
|
err: T,
|
||||||
|
|
|
@ -1681,7 +1681,7 @@ impl<'a, 'b> FlattenArguments<'a, 'b> {
|
||||||
self.func.push(Statement::Constant(ConstantDefinition {
|
self.func.push(Statement::Constant(ConstantDefinition {
|
||||||
dst: id_constant_stmt,
|
dst: id_constant_stmt,
|
||||||
typ: ast::ScalarType::from_parts(width, kind),
|
typ: ast::ScalarType::from_parts(width, kind),
|
||||||
value: -(offset as i64),
|
value: ast::ImmediateValue::S64(-(offset as i64)),
|
||||||
}));
|
}));
|
||||||
self.func.push(Statement::Instruction(
|
self.func.push(Statement::Instruction(
|
||||||
ast::Instruction::<ExpandedArgParams>::Sub(
|
ast::Instruction::<ExpandedArgParams>::Sub(
|
||||||
|
@ -1697,7 +1697,7 @@ impl<'a, 'b> FlattenArguments<'a, 'b> {
|
||||||
self.func.push(Statement::Constant(ConstantDefinition {
|
self.func.push(Statement::Constant(ConstantDefinition {
|
||||||
dst: id_constant_stmt,
|
dst: id_constant_stmt,
|
||||||
typ: ast::ScalarType::from_parts(width, kind),
|
typ: ast::ScalarType::from_parts(width, kind),
|
||||||
value: offset as i64,
|
value: ast::ImmediateValue::S64(offset as i64),
|
||||||
}));
|
}));
|
||||||
self.func.push(Statement::Instruction(
|
self.func.push(Statement::Instruction(
|
||||||
ast::Instruction::<ExpandedArgParams>::Add(
|
ast::Instruction::<ExpandedArgParams>::Add(
|
||||||
|
@ -1724,7 +1724,7 @@ impl<'a, 'b> FlattenArguments<'a, 'b> {
|
||||||
|
|
||||||
fn immediate(
|
fn immediate(
|
||||||
&mut self,
|
&mut self,
|
||||||
desc: ArgumentDescriptor<u32>,
|
desc: ArgumentDescriptor<ast::ImmediateValue>,
|
||||||
typ: &ast::Type,
|
typ: &ast::Type,
|
||||||
) -> Result<spirv::Word, TranslateError> {
|
) -> Result<spirv::Word, TranslateError> {
|
||||||
let scalar_t = if let ast::Type::Scalar(scalar) = typ {
|
let scalar_t = if let ast::Type::Scalar(scalar) = typ {
|
||||||
|
@ -1736,7 +1736,7 @@ impl<'a, 'b> FlattenArguments<'a, 'b> {
|
||||||
self.func.push(Statement::Constant(ConstantDefinition {
|
self.func.push(Statement::Constant(ConstantDefinition {
|
||||||
dst: id,
|
dst: id,
|
||||||
typ: scalar_t,
|
typ: scalar_t,
|
||||||
value: desc.op as i64,
|
value: desc.op,
|
||||||
}));
|
}));
|
||||||
Ok(id)
|
Ok(id)
|
||||||
}
|
}
|
||||||
|
@ -2081,32 +2081,82 @@ fn emit_function_body_ops(
|
||||||
}
|
}
|
||||||
Statement::Constant(cnst) => {
|
Statement::Constant(cnst) => {
|
||||||
let typ_id = map.get_or_add_scalar(builder, cnst.typ);
|
let typ_id = map.get_or_add_scalar(builder, cnst.typ);
|
||||||
match cnst.typ {
|
match (cnst.typ, cnst.value) {
|
||||||
ast::ScalarType::B8 | ast::ScalarType::U8 => {
|
(ast::ScalarType::B8, ast::ImmediateValue::U64(value))
|
||||||
builder.constant_u32(typ_id, Some(cnst.dst), cnst.value as u8 as u32);
|
| (ast::ScalarType::U8, ast::ImmediateValue::U64(value)) => {
|
||||||
|
builder.constant_u32(typ_id, Some(cnst.dst), value as u8 as u32);
|
||||||
}
|
}
|
||||||
ast::ScalarType::B16 | ast::ScalarType::U16 => {
|
(ast::ScalarType::B16, ast::ImmediateValue::U64(value))
|
||||||
builder.constant_u32(typ_id, Some(cnst.dst), cnst.value as u16 as u32);
|
| (ast::ScalarType::U16, ast::ImmediateValue::U64(value)) => {
|
||||||
|
builder.constant_u32(typ_id, Some(cnst.dst), value as u16 as u32);
|
||||||
}
|
}
|
||||||
ast::ScalarType::B32 | ast::ScalarType::U32 => {
|
(ast::ScalarType::B32, ast::ImmediateValue::U64(value))
|
||||||
builder.constant_u32(typ_id, Some(cnst.dst), cnst.value as u32);
|
| (ast::ScalarType::U32, ast::ImmediateValue::U64(value)) => {
|
||||||
|
builder.constant_u32(typ_id, Some(cnst.dst), value as u32);
|
||||||
}
|
}
|
||||||
ast::ScalarType::B64 | ast::ScalarType::U64 => {
|
(ast::ScalarType::B64, ast::ImmediateValue::U64(value))
|
||||||
builder.constant_u64(typ_id, Some(cnst.dst), cnst.value as u64);
|
| (ast::ScalarType::U64, ast::ImmediateValue::U64(value)) => {
|
||||||
|
builder.constant_u64(typ_id, Some(cnst.dst), value);
|
||||||
}
|
}
|
||||||
ast::ScalarType::S8 => {
|
(ast::ScalarType::S8, ast::ImmediateValue::U64(value)) => {
|
||||||
builder.constant_u32(typ_id, Some(cnst.dst), cnst.value as i8 as u32);
|
builder.constant_u32(typ_id, Some(cnst.dst), value as i8 as u32);
|
||||||
}
|
}
|
||||||
ast::ScalarType::S16 => {
|
(ast::ScalarType::S16, ast::ImmediateValue::U64(value)) => {
|
||||||
builder.constant_u32(typ_id, Some(cnst.dst), cnst.value as i16 as u32);
|
builder.constant_u32(typ_id, Some(cnst.dst), value as i16 as u32);
|
||||||
}
|
}
|
||||||
ast::ScalarType::S32 => {
|
(ast::ScalarType::S32, ast::ImmediateValue::U64(value)) => {
|
||||||
builder.constant_u32(typ_id, Some(cnst.dst), cnst.value as i32 as u32);
|
builder.constant_u32(typ_id, Some(cnst.dst), value as i32 as u32);
|
||||||
}
|
}
|
||||||
ast::ScalarType::S64 => {
|
(ast::ScalarType::S64, ast::ImmediateValue::U64(value)) => {
|
||||||
builder.constant_u64(typ_id, Some(cnst.dst), cnst.value as i64 as u64);
|
builder.constant_u64(typ_id, Some(cnst.dst), value as i64 as u64);
|
||||||
}
|
}
|
||||||
_ => unreachable!(),
|
(ast::ScalarType::B8, ast::ImmediateValue::S64(value))
|
||||||
|
| (ast::ScalarType::U8, ast::ImmediateValue::S64(value)) => {
|
||||||
|
builder.constant_u32(typ_id, Some(cnst.dst), value as u8 as u32);
|
||||||
|
}
|
||||||
|
(ast::ScalarType::B16, ast::ImmediateValue::S64(value))
|
||||||
|
| (ast::ScalarType::U16, ast::ImmediateValue::S64(value)) => {
|
||||||
|
builder.constant_u32(typ_id, Some(cnst.dst), value as u16 as u32);
|
||||||
|
}
|
||||||
|
(ast::ScalarType::B32, ast::ImmediateValue::S64(value))
|
||||||
|
| (ast::ScalarType::U32, ast::ImmediateValue::S64(value)) => {
|
||||||
|
builder.constant_u32(typ_id, Some(cnst.dst), value as u32);
|
||||||
|
}
|
||||||
|
(ast::ScalarType::B64, ast::ImmediateValue::S64(value))
|
||||||
|
| (ast::ScalarType::U64, ast::ImmediateValue::S64(value)) => {
|
||||||
|
builder.constant_u64(typ_id, Some(cnst.dst), value as u64);
|
||||||
|
}
|
||||||
|
(ast::ScalarType::S8, ast::ImmediateValue::S64(value)) => {
|
||||||
|
builder.constant_u32(typ_id, Some(cnst.dst), value as i8 as u32);
|
||||||
|
}
|
||||||
|
(ast::ScalarType::S16, ast::ImmediateValue::S64(value)) => {
|
||||||
|
builder.constant_u32(typ_id, Some(cnst.dst), value as i16 as u32);
|
||||||
|
}
|
||||||
|
(ast::ScalarType::S32, ast::ImmediateValue::S64(value)) => {
|
||||||
|
builder.constant_u32(typ_id, Some(cnst.dst), value as i32 as u32);
|
||||||
|
}
|
||||||
|
(ast::ScalarType::S64, ast::ImmediateValue::S64(value)) => {
|
||||||
|
builder.constant_u64(typ_id, Some(cnst.dst), value as u64);
|
||||||
|
}
|
||||||
|
(ast::ScalarType::F16, ast::ImmediateValue::F32(value)) => {
|
||||||
|
builder.constant_f32(typ_id, Some(cnst.dst), f16::from_f32(value).to_f32());
|
||||||
|
}
|
||||||
|
(ast::ScalarType::F32, ast::ImmediateValue::F32(value)) => {
|
||||||
|
builder.constant_f32(typ_id, Some(cnst.dst), value);
|
||||||
|
}
|
||||||
|
(ast::ScalarType::F64, ast::ImmediateValue::F32(value)) => {
|
||||||
|
builder.constant_f64(typ_id, Some(cnst.dst), value as f64);
|
||||||
|
}
|
||||||
|
(ast::ScalarType::F16, ast::ImmediateValue::F64(value)) => {
|
||||||
|
builder.constant_f32(typ_id, Some(cnst.dst), f16::from_f64(value).to_f32());
|
||||||
|
}
|
||||||
|
(ast::ScalarType::F32, ast::ImmediateValue::F64(value)) => {
|
||||||
|
builder.constant_f32(typ_id, Some(cnst.dst), value as f32);
|
||||||
|
}
|
||||||
|
(ast::ScalarType::F64, ast::ImmediateValue::F64(value)) => {
|
||||||
|
builder.constant_f64(typ_id, Some(cnst.dst), value);
|
||||||
|
}
|
||||||
|
_ => return Err(TranslateError::MismatchedType),
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
Statement::Conversion(cv) => emit_implicit_conversion(builder, map, cv)?,
|
Statement::Conversion(cv) => emit_implicit_conversion(builder, map, cv)?,
|
||||||
|
@ -4371,7 +4421,7 @@ impl VisitVariableExpanded for CompositeRead {
|
||||||
struct ConstantDefinition {
|
struct ConstantDefinition {
|
||||||
pub dst: spirv::Word,
|
pub dst: spirv::Word,
|
||||||
pub typ: ast::ScalarType,
|
pub typ: ast::ScalarType,
|
||||||
pub value: i64,
|
pub value: ast::ImmediateValue,
|
||||||
}
|
}
|
||||||
|
|
||||||
struct BrachCondition {
|
struct BrachCondition {
|
||||||
|
|
Loading…
Reference in a new issue