Implement neg instruction

6 files changed, 131 insertions, 1 deletions

diff --git a/ptx/src/ast.rs b/ptx/src/ast.rs
index f00ddce..7f2fc9a 100644
--- a/ptx/src/ast.rs
+++ b/ptx/src/ast.rs
@@ -542,6 +542,7 @@ pub enum Instruction<P: ArgParams> {
     Div(DivDetails, Arg3<P>),
     Sqrt(SqrtDetails, Arg2<P>),
     Rsqrt(RsqrtDetails, Arg2<P>),
+    Neg(NegDetails, Arg2<P>),
 }
 
 #[derive(Copy, Clone)]
@@ -1183,6 +1184,12 @@ pub struct RsqrtDetails {
     pub flush_to_zero: bool,
 }
 
+#[derive(Copy, Clone, Eq, PartialEq)]
+pub struct NegDetails {
+    pub typ: ScalarType,
+    pub flush_to_zero: Option<bool>,
+}
+
 impl<'a> NumsOrArrays<'a> {
     pub fn to_vec(self, typ: SizedScalarType, dimensions: &mut [u32]) -> Result<Vec<u8>, PtxError> {
         self.normalize_dimensions(dimensions)?;
diff --git a/ptx/src/ptx.lalrpop b/ptx/src/ptx.lalrpop
index 4cf4255..9d2adec 100644
--- a/ptx/src/ptx.lalrpop
+++ b/ptx/src/ptx.lalrpop
@@ -156,6 +156,7 @@ match {
     "min",
     "mov",
     "mul",
+    "neg",
     "not",
     "or",
     "rcp",
@@ -198,6 +199,7 @@ ExtendedID : &'input str = {
     "min",
     "mov",
     "mul",
+    "neg",
     "not",
     "or",
     "rcp",
@@ -684,6 +686,7 @@ Instruction: ast::Instruction<ast::ParsedArgParams<'input>> = {
     InstDiv,
     InstSqrt,
     InstRsqrt,
+    InstNeg,
 };
 
 // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-ld
@@ -1577,6 +1580,39 @@ InstRsqrt: ast::Instruction<ast::ParsedArgParams<'input>> = {
     },
 }
 
+// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-neg
+// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-neg
+// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#half-precision-floating-point-instructions-neg
+InstNeg: ast::Instruction<ast::ParsedArgParams<'input>> = {
+    "neg" <ftz:".ftz"?> <typ:NegTypeFtz> <a:Arg2> => {
+        let details = ast::NegDetails {
+            typ,
+            flush_to_zero: Some(ftz.is_some()),
+        };
+        ast::Instruction::Neg(details, a)
+    },
+    "neg" <typ:NegTypeNonFtz> <a:Arg2> => {
+        let details = ast::NegDetails {
+            typ,
+            flush_to_zero: None,
+        };
+        ast::Instruction::Neg(details, a)
+    },
+}
+
+NegTypeFtz: ast::ScalarType = {
+    ".f16" => ast::ScalarType::F16,
+    ".f16x2" => ast::ScalarType::F16x2,
+    ".f32" => ast::ScalarType::F32,
+}
+
+NegTypeNonFtz: ast::ScalarType = {
+    ".s16" => ast::ScalarType::S16,
+    ".s32" => ast::ScalarType::S32,
+    ".s64" => ast::ScalarType::S64,
+    ".f64" => ast::ScalarType::F64
+}
+
 ArithDetails: ast::ArithDetails = {
     <t:UIntType> => ast::ArithDetails::Unsigned(t),
     <t:SIntType> => ast::ArithDetails::Signed(ast::ArithSInt {
diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs
index 4e9d39f..7ba3c4d 100644
--- a/ptx/src/test/spirv_run/mod.rs
+++ b/ptx/src/test/spirv_run/mod.rs
@@ -104,6 +104,7 @@ test_ptx!(atom_add, [2u32, 4u32], [2u32, 6u32]);
 test_ptx!(div_approx, [1f32, 2f32], [0.5f32]);

 test_ptx!(sqrt, [0.25f32], [0.5f32]);

 test_ptx!(rsqrt, [0.25f64], [2f64]);

+test_ptx!(neg, [181i32], [-181i32]);

 

 struct DisplayError<T: Debug> {

     err: T,

diff --git a/ptx/src/test/spirv_run/neg.ptx b/ptx/src/test/spirv_run/neg.ptx
new file mode 100644
index 0000000..60fe162
--- /dev/null
+++ b/ptx/src/test/spirv_run/neg.ptx
@@ -0,0 +1,21 @@
+.version 6.5

+.target sm_30

+.address_size 64

+

+.visible .entry neg(

+	.param .u64 input,

+	.param .u64 output

+)

+{

+	.reg .u64 	        in_addr;

+    .reg .u64 	        out_addr;

+    .reg .s32 	        temp1;

+

+	ld.param.u64 	    in_addr, [input];

+    ld.param.u64 	    out_addr, [output];

+

+    ld.s32              temp1, [in_addr];

+	neg.s32			    temp1, temp1;

+    st.s32              [out_addr], temp1;

+	ret;

+}

diff --git a/ptx/src/test/spirv_run/neg.spvtxt b/ptx/src/test/spirv_run/neg.spvtxt
new file mode 100644
index 0000000..b358858
--- /dev/null
+++ b/ptx/src/test/spirv_run/neg.spvtxt
@@ -0,0 +1,47 @@
+               OpCapability GenericPointer
+               OpCapability Linkage
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int64
+               OpCapability Int8
+         %26 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %1 "not"
+       %void = OpTypeVoid
+      %ulong = OpTypeInt 64 0
+         %29 = OpTypeFunction %void %ulong %ulong
+%_ptr_Function_ulong = OpTypePointer Function %ulong
+%_ptr_Generic_ulong = OpTypePointer Generic %ulong
+          %1 = OpFunction %void None %29
+          %8 = OpFunctionParameter %ulong
+          %9 = OpFunctionParameter %ulong
+         %24 = 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_ulong Function
+          %7 = OpVariable %_ptr_Function_ulong Function
+               OpStore %2 %8
+               OpStore %3 %9
+         %11 = OpLoad %ulong %2
+         %10 = OpCopyObject %ulong %11
+               OpStore %4 %10
+         %13 = OpLoad %ulong %3
+         %12 = OpCopyObject %ulong %13
+               OpStore %5 %12
+         %15 = OpLoad %ulong %4
+         %20 = OpConvertUToPtr %_ptr_Generic_ulong %15
+         %14 = OpLoad %ulong %20
+               OpStore %6 %14
+         %17 = OpLoad %ulong %6
+         %22 = OpCopyObject %ulong %17
+         %21 = OpNot %ulong %22
+         %16 = OpCopyObject %ulong %21
+               OpStore %7 %16
+         %18 = OpLoad %ulong %5
+         %19 = OpLoad %ulong %7
+         %23 = OpConvertUToPtr %_ptr_Generic_ulong %18
+               OpStore %23 %19
+               OpReturn
+               OpFunctionEnd
diff --git a/ptx/src/translate.rs b/ptx/src/translate.rs
index c351ccd..36e15f9 100644
--- a/ptx/src/translate.rs
+++ b/ptx/src/translate.rs
@@ -1511,6 +1511,9 @@ fn convert_to_typed_statements(
                 ast::Instruction::Rsqrt(d, a) => {

                     result.push(Statement::Instruction(ast::Instruction::Rsqrt(d, a.cast())))

                 }

+                ast::Instruction::Neg(d, a) => {

+                    result.push(Statement::Instruction(ast::Instruction::Neg(d, a.cast())))

+                }

             },

             Statement::Label(i) => result.push(Statement::Label(i)),

             Statement::Variable(v) => result.push(Statement::Variable(v)),

@@ -2805,6 +2808,15 @@ fn emit_function_body_ops(
                         &[a.src],

                     )?;

                 }

+                ast::Instruction::Neg(details, arg) => {

+                    let result_type = map.get_or_add_scalar(builder, details.typ);

+                    let negate_func = if details.typ.kind() == ScalarKind::Float {

+                        dr::Builder::f_negate

+                    } else {

+                        dr::Builder::s_negate

+                    };

+                    negate_func(builder, result_type, Some(arg.dst), arg.src)?;

+                }

             },

             Statement::LoadVar(arg, typ) => {

                 let type_id = map.get_or_add(builder, SpirvType::from(typ.clone()));

@@ -3406,7 +3418,7 @@ fn emit_setp(
         (ast::SetpCompareOp::NanGreaterOrEq, _) => {

             builder.f_unord_greater_than_equal(result_type, result_id, operand_1, operand_2)

         }

-        _ => todo!()

+        _ => todo!(),

     }?;

     Ok(())

 }

@@ -4678,6 +4690,9 @@ impl<T: ArgParamsEx> ast::Instruction<T> {
             ast::Instruction::Rsqrt(d, a) => {

                 ast::Instruction::Rsqrt(d, a.map(visitor, &ast::Type::Scalar(d.typ.into()))?)

             }

+            ast::Instruction::Neg(d, a) => {

+                ast::Instruction::Neg(d, a.map(visitor, &ast::Type::Scalar(d.typ))?)

+            }

         })

     }

 }

@@ -4984,6 +4999,9 @@ impl ast::Instruction<ExpandedArgParams> {
                 details.flush_to_zero,

                 ast::ScalarType::from(details.typ).size_of(),

             )),

+            ast::Instruction::Neg(details, _) => details

+                .flush_to_zero

+                .map(|ftz| (ftz, details.typ.size_of())),

         }

     }

 }