Skip to content

GitLab

Switch branch/tag
Find file Normal viewHistoryPermalink
test_cub.py 4.33 KB
Newer Older
root's avatar
clone code from github  
root committed
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128 
import cupy
from cupy import testing
from cupy_backends.cuda.api import runtime
from cupyx import jit


class TestCubWarpReduce:

    @testing.for_all_dtypes(no_bool=True)
    def test_sum(self, dtype):

        @jit.rawkernel()
        def warp_reduce_sum(x, y):
            WarpReduce = jit.cub.WarpReduce[dtype]
            temp_storage = jit.shared_memory(
                dtype=WarpReduce.TempStorage, size=1)
            i, j = jit.blockIdx.x, jit.threadIdx.x
            value = x[i, j]
            aggregater = WarpReduce(temp_storage[0])
            aggregate = aggregater.Sum(value)
            if j == 0:
                y[i] = aggregate

        warp_size = 64 if runtime.is_hip else 32
        h, w = (32, warp_size)
        x = testing.shaped_random((h, w), dtype=dtype)
        y = testing.shaped_random((h,), dtype=dtype)
        expected = cupy.asnumpy(x).sum(axis=-1, dtype=dtype)

        warp_reduce_sum[h, w](x, y)
        testing.assert_allclose(y, expected, rtol=1e-6)

    @testing.for_all_dtypes(no_bool=True)
    def test_reduce_sum(self, dtype):

        @jit.rawkernel()
        def warp_reduce_sum(x, y):
            WarpReduce = jit.cub.WarpReduce[dtype]
            temp_storage = jit.shared_memory(
                dtype=WarpReduce.TempStorage, size=1)
            i, j = jit.blockIdx.x, jit.threadIdx.x
            value = x[i, j]
            aggregater = WarpReduce(temp_storage[0])
            aggregate = aggregater.Reduce(value, jit.cub.Sum())
            if j == 0:
                y[i] = aggregate

        warp_size = 64 if runtime.is_hip else 32
        h, w = (32, warp_size)
        x = testing.shaped_random((h, w), dtype=dtype)
        y = testing.shaped_random((h,), dtype=dtype)
        expected = cupy.asnumpy(x).sum(axis=-1, dtype=dtype)

        warp_reduce_sum[h, w](x, y)
        testing.assert_allclose(y, expected, rtol=1e-6)

    @testing.for_all_dtypes(no_bool=True)
    def test_reduce_max(self, dtype):

        @jit.rawkernel()
        def warp_reduce_max(x, y):
            WarpReduce = jit.cub.WarpReduce[dtype]
            temp_storage = jit.shared_memory(
                dtype=WarpReduce.TempStorage, size=1)
            i, j = jit.blockIdx.x, jit.threadIdx.x
            value = x[i, j]
            aggregater = WarpReduce(temp_storage[0])
            aggregate = aggregater.Reduce(value, jit.cub.Max())
            if j == 0:
                y[i] = aggregate

        warp_size = 64 if runtime.is_hip else 32
        h, w = (32, warp_size)
        x = testing.shaped_random((h, w), dtype=dtype)
        y = testing.shaped_random((h,), dtype=dtype)
        expected = cupy.asnumpy(x).max(axis=-1)

        warp_reduce_max[h, w](x, y)
        testing.assert_allclose(y, expected, rtol=1e-6)


class TestCubBlockReduce:

    @testing.for_all_dtypes(no_bool=True)
    def test_sum(self, dtype):

        @jit.rawkernel()
        def block_reduce_sum(x, y):
            BlockReduce = jit.cub.BlockReduce[dtype, 256]
            temp_storage = jit.shared_memory(
                dtype=BlockReduce.TempStorage, size=1)
            i, j = jit.blockIdx.x, jit.threadIdx.x
            value = x[i, j]
            aggregater = BlockReduce(temp_storage[0])
            aggregate = aggregater.Sum(value)
            if j == 0:
                y[i] = aggregate

        h, w = (32, 256)
        x = testing.shaped_random((h, w), dtype=dtype)
        y = testing.shaped_random((h,), dtype=dtype)
        expected = cupy.asnumpy(x).sum(axis=-1, dtype=dtype)

        block_reduce_sum[h, w](x, y)
        testing.assert_allclose(y, expected, rtol=1e-6)

    @testing.for_all_dtypes(no_bool=True)
    def test_reduce_min(self, dtype):

        @jit.rawkernel()
        def block_reduce_min(x, y):
            BlockReduce = jit.cub.BlockReduce[dtype, 256]
            temp_storage = jit.shared_memory(
                dtype=BlockReduce.TempStorage, size=1)
            i, j = jit.blockIdx.x, jit.threadIdx.x
            value = x[i, j]
            aggregater = BlockReduce(temp_storage[0])
            aggregate = aggregater.Reduce(value, jit.cub.Min())
            if j == 0:
                y[i] = aggregate

        h, w = (32, 256)
        x = testing.shaped_random((h, w), dtype=dtype)
        y = testing.shaped_random((h,), dtype=dtype)
        expected = cupy.asnumpy(x).min(axis=-1)

        block_reduce_min[h, w](x, y)
        testing.assert_allclose(y, expected, rtol=1e-6)