aboutsummaryrefslogtreecommitdiffhomepage
path: root/ptx
diff options
context:
space:
mode:
Diffstat (limited to 'ptx')
-rw-r--r--ptx/src/ptx.lalrpop6
-rw-r--r--ptx/src/test/spirv_run/const.ptx8
-rw-r--r--ptx/src/test/spirv_run/const.spvtxt123
-rw-r--r--ptx/src/test/spirv_run/cvt_f64_f32.ptx22
-rw-r--r--ptx/src/test/spirv_run/cvt_f64_f32.spvtxt55
-rw-r--r--ptx/src/test/spirv_run/cvt_s16_s8.ptx26
-rw-r--r--ptx/src/test/spirv_run/cvt_s16_s8.spvtxt58
-rw-r--r--ptx/src/test/spirv_run/mod.rs2
-rw-r--r--ptx/src/translate.rs76
9 files changed, 317 insertions, 59 deletions
diff --git a/ptx/src/ptx.lalrpop b/ptx/src/ptx.lalrpop
index abefdf8..0bc7655 100644
--- a/ptx/src/ptx.lalrpop
+++ b/ptx/src/ptx.lalrpop
@@ -14,7 +14,7 @@ extern {
match {
r"\s+" => { },
r"//[^\n\r]*[\n\r]*" => { },
- r"/\*([^\*]*\*+[^\*/])*([^\*]*\*+|[^\*])*\*/" => { },
+ r"/\*[^*]*\*+(?:[^/*][^*]*\*+)*/" => { },
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,
@@ -1143,11 +1143,11 @@ InstCvt: ast::Instruction<ast::ParsedArgParams<'input>> = {
}
), a)
},
- "cvt" <s:".sat"?> ".f64" ".f32" <a:Arg2> => {
+ "cvt" <s:".sat"?> <f:".ftz"?> ".f64" ".f32" <a:Arg2> => {
ast::Instruction::Cvt(ast::CvtDetails::FloatFromFloat(
ast::CvtDesc {
rounding: None,
- flush_to_zero: None,
+ flush_to_zero: Some(f.is_some()),
saturate: s.is_some(),
dst: ast::ScalarType::F64,
src: ast::ScalarType::F32
diff --git a/ptx/src/test/spirv_run/const.ptx b/ptx/src/test/spirv_run/const.ptx
index c22ac2b..90ac09d 100644
--- a/ptx/src/test/spirv_run/const.ptx
+++ b/ptx/src/test/spirv_run/const.ptx
@@ -19,10 +19,10 @@
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];
+ ld.const.b16 temp1, [constparams];
+ ld.const.b16 temp2, [constparams+2];
+ ld.const.b16 temp3, [constparams+4];
+ ld.const.b16 temp4, [constparams+6];
st.u16 [out_addr], temp1;
st.u16 [out_addr+2], temp2;
st.u16 [out_addr+4], temp3;
diff --git a/ptx/src/test/spirv_run/const.spvtxt b/ptx/src/test/spirv_run/const.spvtxt
index 9a7f254..49ed9c3 100644
--- a/ptx/src/test/spirv_run/const.spvtxt
+++ b/ptx/src/test/spirv_run/const.spvtxt
@@ -7,41 +7,106 @@
OpCapability Int64
OpCapability Float16
OpCapability Float64
- %21 = OpExtInstImport "OpenCL.std"
+ %53 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
- OpEntryPoint Kernel %1 "clz"
+ OpEntryPoint Kernel %2 "const" %1
+ OpExecutionMode %2 ContractionOff
+ OpDecorate %1 Alignment 8
%void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %ushort = OpTypeInt 16 0
+ %uint_4 = OpConstant %uint 4
+%_arr_ushort_uint_4 = OpTypeArray %ushort %uint_4
+ %ushort_10 = OpConstant %ushort 10
+ %ushort_20 = OpConstant %ushort 20
+ %ushort_30 = OpConstant %ushort 30
+ %ushort_40 = OpConstant %ushort 40
+ %63 = OpConstantComposite %_arr_ushort_uint_4 %ushort_10 %ushort_20 %ushort_30 %ushort_40
+ %uint_4_0 = OpConstant %uint 4
+%_ptr_UniformConstant__arr_ushort_uint_4 = OpTypePointer UniformConstant %_arr_ushort_uint_4
+ %1 = OpVariable %_ptr_UniformConstant__arr_ushort_uint_4 UniformConstant %63
%ulong = OpTypeInt 64 0
- %24 = OpTypeFunction %void %ulong %ulong
+ %67 = 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
+%_ptr_Function_ushort = OpTypePointer Function %ushort
+%_ptr_UniformConstant_ushort = OpTypePointer UniformConstant %ushort
+ %ulong_2 = OpConstant %ulong 2
+ %uchar = OpTypeInt 8 0
+%_ptr_UniformConstant_uchar = OpTypePointer UniformConstant %uchar
+ %ulong_4 = OpConstant %ulong 4
+ %ulong_6 = OpConstant %ulong 6
+%_ptr_Generic_ushort = OpTypePointer Generic %ushort
+ %ulong_2_0 = OpConstant %ulong 2
+%_ptr_Generic_uchar = OpTypePointer Generic %uchar
+ %ulong_4_0 = OpConstant %ulong 4
+ %ulong_6_0 = OpConstant %ulong 6
+ %2 = OpFunction %void None %67
+ %11 = OpFunctionParameter %ulong
+ %12 = OpFunctionParameter %ulong
+ %51 = OpLabel
%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
+ %6 = OpVariable %_ptr_Function_ulong Function
+ %7 = OpVariable %_ptr_Function_ushort Function
+ %8 = OpVariable %_ptr_Function_ushort Function
+ %9 = OpVariable %_ptr_Function_ushort Function
+ %10 = OpVariable %_ptr_Function_ushort Function
+ OpStore %3 %11
+ OpStore %4 %12
+ %13 = OpLoad %ulong %3 Aligned 8
+ OpStore %5 %13
+ %14 = OpLoad %ulong %4 Aligned 8
+ OpStore %6 %14
+ %39 = OpBitcast %_ptr_UniformConstant_ushort %1
+ %15 = OpLoad %ushort %39 Aligned 2
+ OpStore %7 %15
+ %40 = OpBitcast %_ptr_UniformConstant_ushort %1
+ %73 = OpBitcast %_ptr_UniformConstant_uchar %40
+ %74 = OpInBoundsPtrAccessChain %_ptr_UniformConstant_uchar %73 %ulong_2
+ %28 = OpBitcast %_ptr_UniformConstant_ushort %74
+ %16 = OpLoad %ushort %28 Aligned 2
+ OpStore %8 %16
+ %41 = OpBitcast %_ptr_UniformConstant_ushort %1
+ %75 = OpBitcast %_ptr_UniformConstant_uchar %41
+ %76 = OpInBoundsPtrAccessChain %_ptr_UniformConstant_uchar %75 %ulong_4
+ %30 = OpBitcast %_ptr_UniformConstant_ushort %76
+ %17 = OpLoad %ushort %30 Aligned 2
+ OpStore %9 %17
+ %42 = OpBitcast %_ptr_UniformConstant_ushort %1
+ %77 = OpBitcast %_ptr_UniformConstant_uchar %42
+ %78 = OpInBoundsPtrAccessChain %_ptr_UniformConstant_uchar %77 %ulong_6
+ %32 = OpBitcast %_ptr_UniformConstant_ushort %78
+ %18 = OpLoad %ushort %32 Aligned 2
+ OpStore %10 %18
+ %19 = OpLoad %ulong %6
+ %20 = OpLoad %ushort %7
+ %43 = OpConvertUToPtr %_ptr_Generic_ushort %19
+ %44 = OpCopyObject %ushort %20
+ OpStore %43 %44 Aligned 2
+ %21 = OpLoad %ulong %6
+ %22 = OpLoad %ushort %8
+ %45 = OpConvertUToPtr %_ptr_Generic_ushort %21
+ %81 = OpBitcast %_ptr_Generic_uchar %45
+ %82 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %81 %ulong_2_0
+ %34 = OpBitcast %_ptr_Generic_ushort %82
+ %46 = OpCopyObject %ushort %22
+ OpStore %34 %46 Aligned 2
+ %23 = OpLoad %ulong %6
+ %24 = OpLoad %ushort %9
+ %47 = OpConvertUToPtr %_ptr_Generic_ushort %23
+ %83 = OpBitcast %_ptr_Generic_uchar %47
+ %84 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %83 %ulong_4_0
+ %36 = OpBitcast %_ptr_Generic_ushort %84
+ %48 = OpCopyObject %ushort %24
+ OpStore %36 %48 Aligned 2
+ %25 = OpLoad %ulong %6
+ %26 = OpLoad %ushort %10
+ %49 = OpConvertUToPtr %_ptr_Generic_ushort %25
+ %85 = OpBitcast %_ptr_Generic_uchar %49
+ %86 = OpInBoundsPtrAccessChain %_ptr_Generic_uchar %85 %ulong_6_0
+ %38 = OpBitcast %_ptr_Generic_ushort %86
+ %50 = OpCopyObject %ushort %26
+ OpStore %38 %50 Aligned 2
OpReturn
OpFunctionEnd
diff --git a/ptx/src/test/spirv_run/cvt_f64_f32.ptx b/ptx/src/test/spirv_run/cvt_f64_f32.ptx
new file mode 100644
index 0000000..7aba351
--- /dev/null
+++ b/ptx/src/test/spirv_run/cvt_f64_f32.ptx
@@ -0,0 +1,22 @@
+.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry cvt_f64_f32(
+ .param .u64 input,
+ .param .u64 output
+)
+{
+ .reg .u64 in_addr;
+ .reg .u64 out_addr;
+ .reg .f32 temp_f32;
+ .reg .f64 temp_f64;
+
+ ld.param.u64 in_addr, [input];
+ ld.param.u64 out_addr, [output];
+
+ ld.global.f32 temp_f32, [in_addr];
+ cvt.ftz.f64.f32 temp_f64, temp_f32;
+ st.f64 [out_addr], temp_f64;
+ ret;
+}
diff --git a/ptx/src/test/spirv_run/cvt_f64_f32.spvtxt b/ptx/src/test/spirv_run/cvt_f64_f32.spvtxt
new file mode 100644
index 0000000..907cce4
--- /dev/null
+++ b/ptx/src/test/spirv_run/cvt_f64_f32.spvtxt
@@ -0,0 +1,55 @@
+ OpCapability GenericPointer
+ OpCapability Linkage
+ OpCapability Addresses
+ OpCapability Kernel
+ OpCapability Int8
+ OpCapability Int16
+ OpCapability Int64
+ OpCapability Float16
+ OpCapability Float64
+ %22 = OpExtInstImport "OpenCL.std"
+ OpMemoryModel Physical64 OpenCL
+ OpEntryPoint Kernel %1 "cvt_f64_f32"
+ OpExecutionMode %1 DenormFlushToZero 16
+ OpExecutionMode %1 DenormFlushToZero 32
+ OpExecutionMode %1 DenormFlushToZero 64
+ OpExecutionMode %1 ContractionOff
+ %void = OpTypeVoid
+ %ulong = OpTypeInt 64 0
+ %25 = OpTypeFunction %void %ulong %ulong
+%_ptr_Function_ulong = OpTypePointer Function %ulong
+ %float = OpTypeFloat 32
+%_ptr_Function_float = OpTypePointer Function %float
+ %double = OpTypeFloat 64
+%_ptr_Function_double = OpTypePointer Function %double
+%_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float
+%_ptr_Generic_double = OpTypePointer Generic %double
+ %1 = OpFunction %void None %25
+ %8 = OpFunctionParameter %ulong
+ %9 = OpFunctionParameter %ulong
+ %20 = 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_float Function
+ %7 = OpVariable %_ptr_Function_double Function
+ OpStore %2 %8
+ OpStore %3 %9
+ %10 = OpLoad %ulong %2 Aligned 8
+ OpStore %4 %10
+ %11 = OpLoad %ulong %3 Aligned 8
+ OpStore %5 %11
+ %13 = OpLoad %ulong %4
+ %18 = OpConvertUToPtr %_ptr_CrossWorkgroup_float %13
+ %12 = OpLoad %float %18 Aligned 4
+ OpStore %6 %12
+ %15 = OpLoad %float %6
+ %14 = OpFConvert %double %15
+ OpStore %7 %14
+ %16 = OpLoad %ulong %5
+ %17 = OpLoad %double %7
+ %19 = OpConvertUToPtr %_ptr_Generic_double %16
+ OpStore %19 %17 Aligned 8
+ OpReturn
+ OpFunctionEnd
diff --git a/ptx/src/test/spirv_run/cvt_s16_s8.ptx b/ptx/src/test/spirv_run/cvt_s16_s8.ptx
new file mode 100644
index 0000000..44c0891
--- /dev/null
+++ b/ptx/src/test/spirv_run/cvt_s16_s8.ptx
@@ -0,0 +1,26 @@
+.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry cvt_s16_s8(
+ .param .u64 input,
+ .param .u64 output
+)
+{
+ .reg .u64 in_addr;
+ .reg .u64 out_addr;
+ .reg .b32 temp_16;
+ .reg .b32 temp_8;
+
+ // inline asm
+ /*ptx_texBake_end*/
+ // inline asm
+
+ ld.param.u64 in_addr, [input];
+ ld.param.u64 out_addr, [output];
+
+ ld.global.b32 temp_8, [in_addr];
+ cvt.s16.s8 temp_16, temp_8;
+ st.b32 [out_addr], temp_16;
+ ret;
+}
diff --git a/ptx/src/test/spirv_run/cvt_s16_s8.spvtxt b/ptx/src/test/spirv_run/cvt_s16_s8.spvtxt
new file mode 100644
index 0000000..5f4b050
--- /dev/null
+++ b/ptx/src/test/spirv_run/cvt_s16_s8.spvtxt
@@ -0,0 +1,58 @@
+ OpCapability GenericPointer
+ OpCapability Linkage
+ OpCapability Addresses
+ OpCapability Kernel
+ OpCapability Int8
+ OpCapability Int16
+ OpCapability Int64
+ OpCapability Float16
+ OpCapability Float64
+ %24 = OpExtInstImport "OpenCL.std"
+ OpMemoryModel Physical64 OpenCL
+ OpEntryPoint Kernel %1 "cvt_s16_s8"
+ OpExecutionMode %1 ContractionOff
+ %void = OpTypeVoid
+ %ulong = OpTypeInt 64 0
+ %27 = OpTypeFunction %void %ulong %ulong
+%_ptr_Function_ulong = OpTypePointer Function %ulong
+ %uint = OpTypeInt 32 0
+%_ptr_Function_uint = OpTypePointer Function %uint
+%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
+ %uchar = OpTypeInt 8 0
+ %ushort = OpTypeInt 16 0
+%_ptr_Generic_uint = OpTypePointer Generic %uint
+ %1 = OpFunction %void None %27
+ %8 = OpFunctionParameter %ulong
+ %9 = OpFunctionParameter %ulong
+ %22 = 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
+ %7 = OpVariable %_ptr_Function_uint Function
+ OpStore %2 %8
+ OpStore %3 %9
+ %10 = OpLoad %ulong %2 Aligned 8
+ OpStore %4 %10
+ %11 = OpLoad %ulong %3 Aligned 8
+ OpStore %5 %11
+ %13 = OpLoad %ulong %4
+ %18 = OpConvertUToPtr %_ptr_CrossWorkgroup_uint %13
+ %12 = OpLoad %uint %18 Aligned 4
+ OpStore %7 %12
+ %15 = OpLoad %uint %7
+ %32 = OpBitcast %uint %15
+ %34 = OpUConvert %uchar %32
+ %20 = OpCopyObject %uchar %34
+ %35 = OpBitcast %uchar %20
+ %37 = OpSConvert %ushort %35
+ %19 = OpCopyObject %ushort %37
+ %14 = OpSConvert %uint %19
+ OpStore %6 %14
+ %16 = OpLoad %ulong %5
+ %17 = OpLoad %uint %6
+ %21 = OpConvertUToPtr %_ptr_Generic_uint %16
+ OpStore %21 %17 Aligned 4
+ OpReturn
+ OpFunctionEnd
diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs
index 13cf0f1..97cfbb5 100644
--- a/ptx/src/test/spirv_run/mod.rs
+++ b/ptx/src/test/spirv_run/mod.rs
@@ -204,6 +204,8 @@ 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!(cvt_s16_s8, [0x139231C2u32], [0xFFFFFFC2u32]);
+test_ptx!(cvt_f64_f32, [0.125f32], [0.125f64]);
struct DisplayError<T: Debug> {
err: T,
diff --git a/ptx/src/translate.rs b/ptx/src/translate.rs
index 14e8473..a41179d 100644
--- a/ptx/src/translate.rs
+++ b/ptx/src/translate.rs
@@ -37,6 +37,10 @@ fn error_unreachable() -> TranslateError {
TranslateError::Unreachable
}
+fn error_unknown_symbol() -> TranslateError {
+ TranslateError::UnknownSymbol
+}
+
#[derive(PartialEq, Eq, Hash, Clone)]
enum SpirvType {
Base(SpirvScalarKey),
@@ -3301,7 +3305,7 @@ fn emit_variable<'input>(
}
ast::StateSpace::Global => (true, spirv::StorageClass::CrossWorkgroup),
ast::StateSpace::Shared => (false, spirv::StorageClass::Workgroup),
- ast::StateSpace::Const => todo!(),
+ ast::StateSpace::Const => (false, spirv::StorageClass::UniformConstant),
ast::StateSpace::Generic => todo!(),
ast::StateSpace::Sreg => todo!(),
};
@@ -4168,6 +4172,7 @@ fn normalize_identifiers<'input, 'b>(
match s {
ast::Statement::Label(id) => {
id_defs.add_def(*id, None, false);
+ eprintln!("{}", id);
}
_ => (),
}
@@ -4996,7 +5001,7 @@ impl<'input> GlobalStringIdResolver<'input> {
self.variables
.get(id)
.copied()
- .ok_or(TranslateError::UnknownSymbol)
+ .ok_or_else(error_unknown_symbol)
}
fn current_id(&self) -> spirv::Word {
@@ -5058,7 +5063,7 @@ pub struct GlobalFnDeclResolver<'input, 'a> {
impl<'input, 'a> GlobalFnDeclResolver<'input, 'a> {
fn get_fn_sig_resolver(&self, id: spirv::Word) -> Result<&FnSigMapper<'input>, TranslateError> {
- self.fns.get(&id).ok_or(TranslateError::UnknownSymbol)
+ self.fns.get(&id).ok_or_else(error_unknown_symbol)
}
}
@@ -5099,8 +5104,7 @@ impl<'a, 'b> FnStringIdResolver<'a, 'b> {
match self.global_variables.get(id) {
Some(id) => Ok(*id),
None => {
- let sreg =
- PtxSpecialRegister::try_parse(id).ok_or(TranslateError::UnknownSymbol)?;
+ let sreg = PtxSpecialRegister::try_parse(id).ok_or_else(error_unknown_symbol)?;
Ok(self.special_registers.get_or_add(self.current_id, sreg))
}
}
@@ -5778,25 +5782,13 @@ impl<T: ArgParamsEx> ast::Instruction<T> {
ast::Instruction::Not(t, a.map(visitor, &ast::Type::Scalar(t))?)
}
ast::Instruction::Cvt(d, a) => {
- let (dst_t, src_t) = match &d {
- ast::CvtDetails::FloatFromFloat(desc) => (
- ast::Type::Scalar(desc.dst.into()),
- ast::Type::Scalar(desc.src.into()),
- ),
- ast::CvtDetails::FloatFromInt(desc) => (
- ast::Type::Scalar(desc.dst.into()),
- ast::Type::Scalar(desc.src.into()),
- ),
- ast::CvtDetails::IntFromFloat(desc) => (
- ast::Type::Scalar(desc.dst.into()),
- ast::Type::Scalar(desc.src.into()),
- ),
- ast::CvtDetails::IntFromInt(desc) => (
- ast::Type::Scalar(desc.dst.into()),
- ast::Type::Scalar(desc.src.into()),
- ),
+ let (dst_t, src_t, int_to_int) = match &d {
+ ast::CvtDetails::FloatFromFloat(desc) => ((desc.dst, desc.src, false)),
+ ast::CvtDetails::FloatFromInt(desc) => ((desc.dst, desc.src, false)),
+ ast::CvtDetails::IntFromFloat(desc) => ((desc.dst, desc.src, false)),
+ ast::CvtDetails::IntFromInt(desc) => ((desc.dst, desc.src, true)),
};
- ast::Instruction::Cvt(d, a.map_different_types(visitor, &dst_t, &src_t)?)
+ ast::Instruction::Cvt(d, a.map_cvt(visitor, dst_t, src_t, int_to_int)?)
}
ast::Instruction::Shl(t, a) => {
ast::Instruction::Shl(t, a.map_shift(visitor, &ast::Type::Scalar(t))?)
@@ -6413,6 +6405,44 @@ impl<T: ArgParamsEx> ast::Arg2<T> {
})
}
+ fn map_cvt<U: ArgParamsEx, V: ArgumentMapVisitor<T, U>>(
+ self,
+ visitor: &mut V,
+ dst_t: ast::ScalarType,
+ src_t: ast::ScalarType,
+ is_int_to_int: bool,
+ ) -> Result<ast::Arg2<U>, TranslateError> {
+ let dst = visitor.operand(
+ ArgumentDescriptor {
+ op: self.dst,
+ is_dst: true,
+ is_memory_access: false,
+ non_default_implicit_conversion: if is_int_to_int {
+ Some(should_convert_relaxed_dst_wrapper)
+ } else {
+ None
+ },
+ },
+ &ast::Type::Scalar(dst_t),
+ ast::StateSpace::Reg,
+ )?;
+ let src = visitor.operand(
+ ArgumentDescriptor {
+ op: self.src,
+ is_dst: false,
+ is_memory_access: false,
+ non_default_implicit_conversion: if is_int_to_int {
+ Some(should_convert_relaxed_src_wrapper)
+ } else {
+ None
+ },
+ },
+ &ast::Type::Scalar(src_t),
+ ast::StateSpace::Reg,
+ )?;
+ Ok(ast::Arg2 { dst, src })
+ }
+
fn map_different_types<U: ArgParamsEx, V: ArgumentMapVisitor<T, U>>(
self,
visitor: &mut V,
generated by cgit v1.2.3 (git 2.39.1) at 2024-12-26 14:24:19 +0000