Add test_reduce_api.py::test_reduce_struct_type_minmax (NVIDIA#3938)

* Add test_reduce_api.py::test_reduce_struct_type_minmax

This adds an example of using struct that holds minimum and maximum values
to compute smallest and largest values of an array using a single call
cub::DeviceReduce

It also demonstrates a use of TransformedIterator to map input array
values as v -> MinMax(v,v) to make reduction binary operator commutative,
to render data-parallel algorithm applicable. Since this is done in
registers, each element of input array is only accessed once.

* Refined docstring text

* Change reduce-minmax to compute mininum/maximum absolute values of an array

* Replace embedding_op to transform_op, add comment

Author: oleksandr-pavlyk

python/cuda_parallel/cuda/parallel/experimental/struct.py

@@ -38,15 +38,15 @@ def gpu_struct(this: type) -> Type[GpuStruct]:
 
     Example:
         The code snippet below shows how to use `gpu_struct` to define
-        a `Pixel` type (composed of `r`, `g` and `b` values),  and perform
-        a reduction on an array of `Pixel` objects to identify the one
-        with the largest `g` component:
+        a `MinMax` type (composed of `min_val`, `max_val` values), and perform
+        a reduction on an input array of floating point values to compute its
+        the smallest and the largest absolute values:
 
         .. literalinclude:: ../../python/cuda_parallel/tests/test_reduce_api.py
             :language: python
             :dedent:
-            :start-after: example-begin reduce-struct
-            :end-before: example-end reduce-struct
+            :start-after: example-begin reduce-minmax
+            :end-before: example-end reduce-minmax
     """
     # Implementation-wise, @gpu_struct creates and registers a
     # corresponding numba type to the given type, so that it can be

python/cuda_parallel/tests/test_reduce_api.py

@@ -213,3 +213,61 @@ def max_g_value(x, y):
 
     np.testing.assert_equal(expected["g"], d_out.get()["g"])
     # example-end reduce-struct
+
+
+def test_reduce_struct_type_minmax():
+    # example-begin reduce-minmax
+    import cupy as cp
+    import numpy as np
+
+    import cuda.parallel.experimental.algorithms as algorithms
+    import cuda.parallel.experimental.iterators as iterators
+    from cuda.parallel.experimental.struct import gpu_struct
+
+    @gpu_struct
+    class MinMax:
+        min_val: np.float64
+        max_val: np.float64
+
+    def minmax_op(v1: MinMax, v2: MinMax):
+        c_min = min(v1.min_val, v2.min_val)
+        c_max = max(v1.max_val, v2.max_val)
+        return MinMax(c_min, c_max)
+
+    def transform_op(v):
+        av = abs(v)
+        return MinMax(av, av)
+
+    nelems = 4096
+
+    d_in = cp.random.randn(nelems)
+    # input values must be transformed to MinMax structures
+    # in-place to map computation to data-parallel reduction
+    # algorithm that requires commutative binary operation
+    # with both operands having the same type.
+    tr_it = iterators.TransformIterator(d_in, transform_op)
+
+    d_out = cp.empty(tuple(), dtype=MinMax.dtype)
+
+    # initial value set with identity elements of
+    # minimum and maximum operators
+    h_init = MinMax(np.inf, -np.inf)
+
+    # get algorithm object
+    cccl_sum = algorithms.reduce_into(tr_it, d_out, minmax_op, h_init)
+
+    # allocated needed temporary
+    tmp_sz = cccl_sum(None, tr_it, d_out, nelems, h_init)
+    tmp_storage = cp.empty(tmp_sz, dtype=cp.uint8)
+
+    # invoke the reduction algorithm
+    cccl_sum(tmp_storage, tr_it, d_out, nelems, h_init)
+
+    # display values computed on the device
+    actual = d_out.get()
+
+    h = np.abs(d_in.get())
+    expected = np.asarray([(h.min(), h.max())], dtype=MinMax.dtype)
+
+    assert actual == expected
+    # example-end reduce-minmax