import unittest
import numpy

import cupy
from cupyx import jit
from cupy import testing


class TestRaw(unittest.TestCase):

    def test_raw_kick_one_kernel(self):
        @jit.rawkernel()
        def f(x, y):
            y[0] = x[0]

        x = cupy.array([10], dtype=numpy.int32)
        y = cupy.array([20], dtype=numpy.int32)
        f((1,), (1,), (x, y))
        assert int(y[0]) == 10

    def test_raw_elementwise_single_op(self):
        @jit.rawkernel()
        def f(x, y):
            tid = jit.threadIdx.x + jit.blockDim.x * jit.blockIdx.x
            y[tid] = x[tid]

        x = testing.shaped_random((30,), dtype=numpy.int32, seed=0)
        y = testing.shaped_random((30,), dtype=numpy.int32, seed=1)
        f((5,), (6,), (x, y))
        assert bool((x == y).all())

    def test_raw_elementwise_loop(self):
        @jit.rawkernel()
        def f(x, y, size):
            tid = jit.threadIdx.x + jit.blockDim.x * jit.blockIdx.x
            ntid = jit.blockDim.x * jit.gridDim.x
            for i in range(tid, size, ntid):
                y[i] = x[i]

        x = testing.shaped_random((1024,), dtype=numpy.int32, seed=0)
        y = testing.shaped_random((1024,), dtype=numpy.int32, seed=1)
        f((5,), (6,), (x, y, numpy.uint32(1024)))
        assert bool((x == y).all())

    def test_raw_grid_block_interface(self):
        @jit.rawkernel()
        def f(x, y, size):
            tid = jit.threadIdx.x + jit.blockDim.x * jit.blockIdx.x
            ntid = jit.blockDim.x * jit.gridDim.x
            for i in range(tid, size, ntid):
                y[i] = x[i]

        x = testing.shaped_random((1024,), dtype=numpy.int32, seed=0)
        y = testing.shaped_random((1024,), dtype=numpy.int32, seed=1)
        f[5, 6](x, y, numpy.uint32(1024))
        assert bool((x == y).all())

@@ -140,23 +140,18 @@
 raw_
-kick_one_kernel
+onw_thread
 (sel
@@ -1279,32 +1279,716 @@
 x == y).all())%0A%0A
+    def test_raw_multidimensional_array(self):%0A        @jit.rawkernel()%0A        def f(x, y, n_row, n_col):%0A            tid = jit.threadIdx.x + jit.blockDim.x * jit.blockIdx.x%0A            ntid = jit.blockDim.x * jit.gridDim.x%0A            size = n_row * n_col%0A            for i in range(tid, size, ntid):%0A                i_row = i // n_col%0A                i_col = i %25 n_col%0A                y%5Bi_row, i_col%5D = x%5Bi_row, i_col%5D%0A%0A        n, m = numpy.uint32(12), numpy.uint32(13)%0A        x = testing.shaped_random((n, m), dtype=numpy.int32, seed=0)%0A        y = testing.shaped_random((n, m), dtype=numpy.int32, seed=1)%0A        f((5,), (6,), (x, y, n, m))%0A        assert bool((x == y).all())%0A%0A
     def test_raw

Add test