Add test and full dtype support for argsort, argmax, argmin

cinn/runtime/cuda/cinn_cuda_runtime_source.cuh

@@ -686,6 +686,8 @@ __device__ inline int cinn_cuda_find_float_from(const float *buf, int size, floa
 
 CINN_NVGPU_LT_NUM(fp32, float)
 CINN_NVGPU_LT_NUM(fp64, double)
+CINN_NVGPU_LT_NUM(uint8, uint8_t)
+CINN_NVGPU_LT_NUM(int16, int16_t)
 CINN_NVGPU_LT_NUM(int32, int)
 CINN_NVGPU_LT_NUM(int64, long long int)
 #ifdef CINN_CUDA_FP16
@@ -706,6 +708,8 @@ CINN_NVGPU_LT_NUM(fp16, float16)
 
 CINN_NVGPU_GT_NUM(fp32, float)
 CINN_NVGPU_GT_NUM(fp64, double)
+CINN_NVGPU_GT_NUM(uint8, uint8_t)
+CINN_NVGPU_GT_NUM(int16, int16_t)
 CINN_NVGPU_GT_NUM(int32, int)
 CINN_NVGPU_GT_NUM(int64, long long int)
 #ifdef CINN_CUDA_FP16

cinn/runtime/cuda/cuda_intrinsics.cc

@@ -327,6 +327,8 @@ CINN_REGISTER_HELPER(cuda_intrinsics) {
 
   _REGISTER_CINN_NVGPU_LT_NUM(fp32, float);
   _REGISTER_CINN_NVGPU_LT_NUM(fp64, double);
+  _REGISTER_CINN_NVGPU_LT_NUM(uint8, uint8_t);
+  _REGISTER_CINN_NVGPU_LT_NUM(int16, int16_t);
   _REGISTER_CINN_NVGPU_LT_NUM(int32, int);
   _REGISTER_CINN_NVGPU_LT_NUM(int64, int64_t);
 
@@ -344,6 +346,8 @@ CINN_REGISTER_HELPER(cuda_intrinsics) {
 
   _REGISTER_CINN_NVGPU_GT_NUM(fp32, float);
   _REGISTER_CINN_NVGPU_GT_NUM(fp64, double);
+  _REGISTER_CINN_NVGPU_GT_NUM(uint8, uint8_t);
+  _REGISTER_CINN_NVGPU_GT_NUM(int16, int16_t);
   _REGISTER_CINN_NVGPU_GT_NUM(int32, int);
   _REGISTER_CINN_NVGPU_GT_NUM(int64, int64_t);
 

python/tests/ops/test_argmax_op.py

@@ -0,0 +1,189 @@
+#!/usr/bin/env python3
+
+# Copyright (c) 2023 CINN Authors. All Rights Reserved.
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+import paddle
+import numpy as np
+from op_test import OpTest, OpTestTool
+from op_test_helper import TestCaseHelper
+from cinn.frontend import *
+from cinn.common import *
+
+
+@OpTestTool.skip_if(not is_compiled_with_cuda(),
+                    "x86 test will be skipped due to timeout.")
+class TestArgMaxOp(OpTest):
+    def setUp(self):
+        # print(f"\n{self.__class__.__name__}: {self.case}")
+        self.prepare_inputs()
+
+    def prepare_inputs(self):
+        self.x_np = self.random(
+            self.case["shape"], self.case["dtype"], low=0, high=10)
+        self.axis = self.case["axis"]
+        self.keepdim = self.case["keepdim"]
+
+    def build_paddle_program(self, target):
+        x = paddle.to_tensor(self.x_np, stop_gradient=True)
+        out = paddle.argmax(x, self.axis, self.keepdim)
+        self.paddle_outputs = [out]
+
+    def build_cinn_program(self, target):
+        builder = NetBuilder("argmax")
+        x = builder.create_input(
+            self.nptype2cinntype(self.case["dtype"]), self.case["shape"], "x")
+        out = builder.argmax(x, self.axis, self.keepdim)
+        prog = builder.build()
+        forward_res = self.get_cinn_output(prog, target, [x], [self.x_np],
+                                           [out])
+        self.cinn_outputs = np.array(forward_res).astype("int64")
+
+    def test_check_results(self):
+        self.check_outputs_and_grads()
+
+
+class TestArgMaxOpShapeTest(TestCaseHelper):
+    def init_attrs(self):
+        self.class_name = "ArgMaxOpShapeTest"
+        self.cls = TestArgMaxOp
+        self.inputs = [
+            {
+                "shape": [512],
+            },
+            {
+                "shape": [1024],
+            },
+            {
+                "shape": [1200],
+            },
+            {
+                "shape": [64, 16],
+            },
+            {
+                "shape": [4, 32, 8],
+            },
+            {
+                "shape": [16, 8, 4, 2],
+            },
+            {
+                "shape": [2, 8, 4, 2, 5],
+            },
+            {
+                "shape": [4, 8, 1, 2, 16],
+            },
+            {
+                "shape": [1],
+            },
+            {
+                "shape": [1, 1, 1, 1],
+            },
+            {
+                "shape": [1, 1, 1, 1, 1],
+            },
+        ]
+        self.dtypes = [{"dtype": "float32"}]
+        self.attrs = [{"axis": 0, "keepdim": False}]
+
+
+class TestArgMaxOpDtypeTest(TestCaseHelper):
+    def init_attrs(self):
+        self.class_name = "ArgMaxOpDtypeTest"
+        self.cls = TestArgMaxOp
+        self.inputs = [
+            {
+                "shape": [1024],
+            },
+        ]
+        self.dtypes = [
+            {
+                "dtype": "float16",
+            },
+            {
+                "dtype": "float32",
+            },
+            {
+                "dtype": "float64",
+            },
+            {
+                "dtype": "uint8",
+            },
+            {
+                "dtype": "int16",
+            },
+            {
+                "dtype": "int32",
+            },
+            {
+                "dtype": "int64",
+            },
+        ]
+        self.attrs = [{"axis": 0, "keepdim": False}]
+
+
+class TestArgMaxOpAxisTest(TestCaseHelper):
+    def init_attrs(self):
+        self.class_name = "ArgMaxOpAxisTest"
+        self.cls = TestArgMaxOp
+        self.inputs = [
+            {
+                "shape": [16, 8, 4, 2],
+            },
+        ]
+        self.dtypes = [{"dtype": "float32"}]
+        self.attrs = [{
+            "axis": 0,
+            "keepdim": False
+        }, {
+            "axis": 1,
+            "keepdim": False
+        }, {
+            "axis": 2,
+            "keepdim": False
+        }, {
+            "axis": 3,
+            "keepdim": False
+        }]
+
+
+class TestArgMaxOpKeepdimTest(TestCaseHelper):
+    def init_attrs(self):
+        self.class_name = "ArgMaxOpKeepdimTest"
+        self.cls = TestArgMaxOp
+        self.inputs = [
+            {
+                "shape": [16, 8, 4, 2],
+            },
+        ]
+        self.dtypes = [{"dtype": "float32"}]
+        self.attrs = [{
+            "axis": 0,
+            "keepdim": True
+        }, {
+            "axis": 1,
+            "keepdim": True
+        }, {
+            "axis": 2,
+            "keepdim": True
+        }, {
+            "axis": 3,
+            "keepdim": True
+        }]
+
+
+if __name__ == "__main__":
+    TestArgMaxOpShapeTest().run()
+    TestArgMaxOpDtypeTest().run()
+    TestArgMaxOpAxisTest().run()
+    TestArgMaxOpKeepdimTest().run()