PKU-DAIR
/
Hetu
Not watched
1
0
Fork 0
Code
Releases 0 Wiki evaluate Activity
Issues 0 Pull Requests 0 Datasets Model Cloudbrain HPC
You can not select more than 25 topics Topics must start with a chinese character,a letter or number, can include dashes ('-') and can be up to 35 characters long.
Tree: 311bfde6c1
Branches Tags
${ item.name }
Create branch ${ searchTerm }
from '311bfde6c1'
${ noResults }
Hetu/tests/test_gpu_op.py

1341 lines
50 kB
Raw Blame History 

  1. import numpy as np

  2. import hetu as ht

  3. from hetu import gpu_links as gpu_op

  4. 

  5. 

  6. def test_array_set():

  7.     ctx = ht.gpu(0)

  8.     shape = (500, 200)

  9.     # oneslike

  10.     arr_x = ht.empty(shape, ctx=ctx)

  11.     gpu_op.array_set(arr_x, 1.)

  12.     x = arr_x.asnumpy()

  13.     np.testing.assert_allclose(np.ones(shape), x)

  14.     # zeroslike

  15.     gpu_op.array_set(arr_x, 0.)

  16.     x = arr_x.asnumpy()

  17.     np.testing.assert_allclose(np.zeros(shape), x)

  18. 

  19. 

  20. def test_broadcast_to():

  21.     ctx = ht.gpu(0)

  22.     shape = (200, 300)

  23.     to_shape = (130, 200, 300)

  24.     x = np.random.uniform(-1, 1, shape).astype(np.float32)

  25.     arr_x = ht.array(x, ctx=ctx)

  26.     arr_y = ht.empty(to_shape, ctx=ctx)

  27.     gpu_op.broadcast_to(arr_x, arr_y)

  28.     y = arr_y.asnumpy()

  29.     np.testing.assert_allclose(np.broadcast_to(x, to_shape), y)

  30. 

  31. 

  32. def test_reduce_sum_axis_zero():

  33.     ctx = ht.gpu(0)

  34.     shape = (20, 1, 1)

  35.     temp_shape = list(shape)

  36.     temp_shape[0] = (temp_shape[0] + 1) // 2

  37.     temp_shape = tuple(temp_shape)

  38.     to_shape = (1, 1)

  39.     x = np.random.uniform(0, 20, shape).astype(np.float32)

  40.     arr_x = ht.array(x, ctx=ctx)

  41.     arr_y = ht.empty(to_shape, ctx=ctx)

  42.     arr_workspace = ht.empty(shape=temp_shape, ctx=ctx)

  43.     gpu_op.reduce_sum_axis_zero(arr_x, arr_y, arr_workspace)

  44.     y = arr_y.asnumpy()

  45.     y_ = np.sum(x, axis=0)

  46.     for index, _ in np.ndenumerate(y):

  47.         v = y[index]

  48.         v_ = y_[index]

  49.         if abs((v - v_) / v_) > 1e-4:

  50.             print(index, v, v_)

  51.     np.testing.assert_allclose(np.sum(x, axis=0), y, rtol=1e-5)

  52. 

  53. 

  54. def test_matrix_elementwise_add():

  55.     ctx = ht.gpu(0)

  56.     shape = (500, 200)

  57.     x = np.random.uniform(0, 10, size=shape).astype(np.float32)

  58.     y = np.random.uniform(0, 10, size=shape).astype(np.float32)

  59.     arr_x = ht.array(x, ctx=ctx)

  60.     arr_y = ht.array(y, ctx=ctx)

  61.     arr_z = ht.empty(shape, ctx=ctx)

  62.     gpu_op.matrix_elementwise_add(arr_x, arr_y, arr_z)

  63.     z = arr_z.asnumpy()

  64.     np.testing.assert_allclose(x + y, z, rtol=1e-5)

  65. 

  66. 

  67. def test_matrix_elementwise_add_by_const():

  68.     shape = (2000, 3000)

  69.     ctx = ht.gpu(0)

  70.     x = np.random.uniform(0, 10, size=shape).astype(np.float32)

  71.     val = np.random.uniform(-5, 5)

  72.     arr_x = ht.array(x, ctx=ctx)

  73.     arr_y = ht.empty(shape, ctx=ctx)

  74.     gpu_op.matrix_elementwise_add_by_const(arr_x, val, arr_y)

  75.     y = arr_y.asnumpy()

  76.     np.testing.assert_allclose(x + val, y, rtol=1e-5)

  77. 

  78. 

  79. def test_matrix_elementwise_multiply():

  80.     ctx = ht.gpu(0)

  81.     shape = (500, 200)

  82.     x = np.random.uniform(0, 10, size=shape).astype(np.float32)

  83.     y = np.random.uniform(0, 10, size=shape).astype(np.float32)

  84.     arr_x = ht.array(x, ctx=ctx)

  85.     arr_y = ht.array(y, ctx=ctx)

  86.     arr_z = ht.empty(shape, ctx=ctx)

  87.     gpu_op.matrix_elementwise_multiply(arr_x, arr_y, arr_z)

  88.     z = arr_z.asnumpy()

  89.     np.testing.assert_allclose(x * y, z, rtol=1e-5)

  90. 

  91. 

  92. def test_matrix_elementwise_multiply_by_const():

  93.     shape = (2000, 3000)

  94.     ctx = ht.gpu(0)

  95.     x = np.random.uniform(0, 10, size=shape).astype(np.float32)

  96.     val = np.random.uniform(-5, 5)

  97.     arr_x = ht.array(x, ctx=ctx)

  98.     arr_y = ht.empty(shape, ctx=ctx)

  99.     gpu_op.matrix_elementwise_multiply_by_const(arr_x, val, arr_y)

  100.     y = arr_y.asnumpy()

  101.     np.testing.assert_allclose(x * val, y, rtol=1e-5)

  102. 

  103. 

  104. def test_matrix_elementwise_divide():

  105.     ctx = ht.gpu(0)

  106.     shape = (500, 200)

  107.     x = np.random.uniform(0, 10, size=shape).astype(np.float32)

  108.     y = np.random.uniform(1, 10, size=shape).astype(np.float32)

  109.     arr_x = ht.array(x, ctx=ctx)

  110.     arr_y = ht.array(y, ctx=ctx)

  111.     arr_z = ht.empty(shape, ctx=ctx)

  112.     gpu_op.matrix_elementwise_divide(arr_x, arr_y, arr_z)

  113.     z = arr_z.asnumpy()

  114.     np.testing.assert_allclose(x / y, z, rtol=1e-5)

  115. 

  116. 

  117. def test_matrix_elementwise_divide_const():

  118.     shape = (2000, 3000)

  119.     ctx = ht.gpu(0)

  120.     val = np.random.uniform(-5, 5)

  121.     x = np.random.uniform(1, 10, size=shape).astype(np.float32)

  122.     arr_x = ht.array(x, ctx=ctx)

  123.     arr_y = ht.empty(shape, ctx=ctx)

  124.     gpu_op.matrix_elementwise_divide_const(val, arr_x, arr_y)

  125.     y = arr_y.asnumpy()

  126.     np.testing.assert_allclose(val / x, y, rtol=1e-5)

  127. 

  128. 

  129. def test_matrix_opposite():

  130.     shape = (2000, 2500)

  131.     ctx = ht.gpu(0)

  132.     x = np.random.uniform(-1, 1, shape).astype(np.float32)

  133.     arr_x = ht.array(x, ctx=ctx)

  134.     arr_y = ht.empty(shape, ctx=ctx)

  135.     gpu_op.matrix_opposite(arr_x, arr_y)

  136.     y = arr_y.asnumpy()

  137.     np.testing.assert_allclose(-x, y)

  138. 

  139. 

  140. def test_matrix_multiply():

  141.     ctx = ht.gpu(0)

  142.     x = np.random.uniform(0, 10, size=(500, 700)).astype(np.float32)

  143.     y = np.random.uniform(0, 10, size=(700, 1000)).astype(np.float32)

  144.     arr_x = ht.array(x, ctx=ctx)

  145.     arr_y = ht.array(y, ctx=ctx)

  146.     arr_z = ht.empty((500, 1000), ctx=ctx)

  147.     gpu_op.matrix_multiply(arr_x, False, arr_y, False, arr_z)

  148.     z = arr_z.asnumpy()

  149.     np.testing.assert_allclose(np.dot(x, y), z, rtol=1e-5)

  150. 

  151.     x = np.random.uniform(0, 10, size=(1000, 500)).astype(np.float32)

  152.     y = np.random.uniform(0, 10, size=(2000, 500)).astype(np.float32)

  153.     arr_x = ht.array(x, ctx=ctx)

  154.     arr_y = ht.array(y, ctx=ctx)

  155.     arr_z = ht.empty((1000, 2000), ctx=ctx)

  156.     gpu_op.matrix_multiply(arr_x, False, arr_y, True, arr_z)

  157.     z = arr_z.asnumpy()

  158.     np.testing.assert_allclose(np.dot(x, np.transpose(y)), z, rtol=1e-5)

  159. 

  160.     x = np.random.uniform(0, 10, size=(500, 1000)).astype(np.float32)

  161.     y = np.random.uniform(0, 10, size=(2000, 500)).astype(np.float32)

  162.     arr_x = ht.array(x, ctx=ctx)

  163.     arr_y = ht.array(y, ctx=ctx)

  164.     arr_z = ht.empty((1000, 2000), ctx=ctx)

  165.     gpu_op.matrix_multiply(arr_x, True, arr_y, True, arr_z)

  166.     z = arr_z.asnumpy()

  167.     np.testing.assert_allclose(np.dot(np.transpose(x), np.transpose(y)), z,

  168.                                rtol=1e-5)

  169. 

  170. 

  171. def test_matrix_sqrt():

  172.     shape = (2000, 2500)

  173.     ctx = ht.gpu(0)

  174.     x = np.random.uniform(0, 10, shape).astype(np.float32)

  175.     arr_x = ht.array(x, ctx=ctx)

  176.     arr_y = ht.empty(shape, ctx=ctx)

  177.     gpu_op.matrix_sqrt(arr_x, arr_y)

  178.     y = arr_y.asnumpy()

  179.     np.testing.assert_allclose(np.sqrt(x), y, rtol=1e-5)

  180. 

  181. 

  182. def test_matrix_rsqrt():

  183.     shape = (2000, 2500)

  184.     ctx = ht.gpu(0)

  185.     x = np.random.uniform(0, 10, shape).astype(np.float32)

  186.     arr_x = ht.array(x, ctx=ctx)

  187.     arr_y = ht.empty(shape, ctx=ctx)

  188.     gpu_op.matrix_rsqrt(arr_x, arr_y)

  189.     y = arr_y.asnumpy()

  190.     np.testing.assert_allclose(1 / np.sqrt(x), y, rtol=1e-5)

  191. 

  192. 

  193. def test_relu():

  194.     shape = (2000, 2500)

  195.     ctx = ht.gpu(0)

  196.     x = np.random.uniform(-1, 1, shape).astype(np.float32)

  197.     arr_x = ht.array(x, ctx=ctx)

  198.     arr_y = ht.empty(shape, ctx=ctx)

  199.     gpu_op.relu(arr_x, arr_y)

  200.     y = arr_y.asnumpy()

  201.     np.testing.assert_allclose(np.maximum(x, 0).astype(np.float32), y)

  202. 

  203. 

  204. def test_leaky_relu():

  205.     shape = (2000, 2500)

  206.     ctx = ht.gpu(0)

  207.     x = np.random.uniform(-1, 1, shape).astype(np.float32)

  208.     arr_x = ht.array(x, ctx=ctx)

  209.     arr_y = ht.empty(shape, ctx=ctx)

  210.     alpha = 10

  211.     gpu_op.leaky_relu(arr_x, float(alpha), arr_y)

  212.     y = arr_y.asnumpy()

  213. 

  214. 

  215. def test_relu_gradient():

  216.     shape = (2000, 2500)

  217.     ctx = ht.gpu(0)

  218.     x = np.random.uniform(-1, 1, shape).astype(np.float32)

  219.     grad_x = np.random.uniform(-5, 5, shape).astype(np.float32)

  220.     arr_x = ht.array(x, ctx=ctx)

  221.     arr_grad_x = ht.array(grad_x, ctx=ctx)

  222.     arr_y = ht.empty(shape, ctx=ctx)

  223.     gpu_op.relu_gradient(arr_x, arr_grad_x, arr_y)

  224.     y = arr_y.asnumpy()

  225.     np.testing.assert_allclose(((x > 0) * grad_x).astype(np.float32), y)

  226. 

  227. 

  228. def test_leaky_relu_gradient():

  229.     shape = (2000, 2500)

  230.     ctx = ht.gpu(0)

  231.     x = np.random.uniform(-1, 1, shape).astype(np.float32)

  232.     grad_x = np.random.uniform(-5, 5, shape).astype(np.float32)

  233.     arr_x = ht.array(x, ctx=ctx)

  234.     arr_grad_x = ht.array(grad_x, ctx=ctx)

  235.     arr_y = ht.empty(shape, ctx=ctx)

  236.     alpha = 10

  237.     gpu_op.leaky_relu_gradient(arr_x, arr_grad_x, alpha, arr_y)

  238.     y = arr_y.asnumpy()

  239.     np.testing.assert_allclose(

  240.         np.where(np.greater(x, 0), grad_x, alpha * grad_x).astype(np.float32), y)

  241. 

  242. 

  243. def test_softmax():

  244.     def softmax_func(y):

  245.         """Numerically stable softmax."""

  246.         b = y - np.max(y, axis=1, keepdims=True)

  247.         expb = np.exp(b)

  248.         softmax = expb / np.sum(expb, axis=1, keepdims=True)

  249.         return softmax

  250.     ctx = ht.gpu(0)

  251.     shape = (400, 1000)

  252.     x = np.random.uniform(-5, 5, shape).astype(np.float32)

  253.     arr_x = ht.array(x, ctx=ctx)

  254.     arr_y = ht.empty(shape, ctx=ctx)

  255.     gpu_op.softmax(arr_x, arr_y)

  256.     y = arr_y.asnumpy()

  257.     np.testing.assert_allclose(softmax_func(x), y, rtol=1e-5)

  258. 

  259. 

  260. def test_softmax_cross_entropy():

  261.     def softmax_func(y):

  262.         """Numerically stable softmax."""

  263.         b = y - np.max(y, axis=1, keepdims=True)

  264.         expb = np.exp(b)

  265.         softmax = expb / np.sum(expb, axis=1, keepdims=True)

  266.         return softmax

  267.     ctx = ht.gpu(0)

  268.     shape = (400, 1000)

  269.     y = np.random.uniform(-5, 5, shape).astype(np.float32)

  270.     y_ = np.random.uniform(-5, 5, shape).astype(np.float32)

  271.     arr_y = ht.array(y, ctx=ctx)

  272.     arr_y_ = ht.array(y_, ctx=ctx)

  273.     arr_out = ht.empty((400,), ctx=ctx)

  274.     gpu_op.softmax_cross_entropy(arr_y, arr_y_, arr_out)

  275.     out = arr_out.asnumpy()

  276.     # numpy calculation

  277.     cross_entropy = -np.sum(y_ * np.log(softmax_func(y)), axis=1)

  278.     np.testing.assert_allclose(cross_entropy, out, rtol=1e-4)

  279.     # test cudnn

  280.     gpu_op.CuDNN_softmax_cross_entropy(arr_y, arr_y_, arr_out)

  281.     out = arr_out.asnumpy()

  282.     np.testing.assert_allclose(cross_entropy, out, rtol=1e-4)

  283. 

  284. 

  285. def test_softmax_cross_entropy_gradient():

  286.     def softmax_func(y):

  287.         """Numerically stable softmax."""

  288.         b = y - np.max(y, axis=1, keepdims=True)

  289.         expb = np.exp(b)

  290.         softmax = expb / np.sum(expb, axis=1, keepdims=True)

  291.         return softmax

  292.     ctx = ht.gpu(0)

  293.     shape = (400, 1000)

  294.     y = np.random.uniform(-5, 5, shape).astype(np.float32)

  295.     y_ = np.random.uniform(-5, 5, shape).astype(np.float32)

  296.     grad = np.random.uniform(-5, 5, (400,)).astype(np.float32)

  297.     arr_y = ht.array(y, ctx=ctx)

  298.     arr_y_ = ht.array(y_, ctx=ctx)

  299.     arr_grad = ht.array(grad, ctx=ctx)

  300.     arr_out = ht.empty(shape, ctx=ctx)

  301.     gpu_op.softmax_cross_entropy_gradient(arr_y, arr_y_, arr_grad, arr_out)

  302.     out = arr_out.asnumpy()

  303.     # numpy calculation

  304.     np_grad = (softmax_func(y) + -1 * y_) * np.expand_dims(grad, -1)

  305.     np.testing.assert_allclose(np_grad, out, rtol=1e-4, atol=1e-8)

  306. 

  307.     # test cudnn

  308.     gpu_op.CuDNN_softmax_cross_entropy_gradient(

  309.         arr_grad, arr_y, arr_y_, arr_out)

  310.     out = arr_out.asnumpy()

  311.     np.testing.assert_allclose(np_grad, out, rtol=1e-4, atol=1e-8)

  312. 

  313. 

  314. def test_conv2d():

  315.     ctx = ht.gpu(0)

  316.     # im2col and np_conv2d are helper functions

  317. 

  318.     def im2col(X, filter_H, filter_W, padding, stride):

  319.         N, C, H, W = X.shape

  320.         assert (H + 2 * padding - filter_H) % stride == 0

  321.         assert (W + 2 * padding - filter_W) % stride == 0

  322.         out_H = (H + 2 * padding - filter_H) // stride + 1

  323.         out_W = (W + 2 * padding - filter_W) // stride + 1

  324. 

  325.         y_row_size = C * filter_H * filter_W

  326.         y_col_size = out_H * out_W

  327.         y_shape = (N, y_row_size, y_col_size)

  328.         Y = np.empty(y_shape, dtype=X.dtype)

  329. 

  330.         for batch_index in range(N):

  331.             for col_index in range(y_col_size):

  332.                 out_y = col_index // out_W

  333.                 out_x = col_index % out_W

  334.                 in_y = out_y * stride - padding

  335.                 in_x = out_x * stride - padding

  336.                 row_idx = 0

  337.                 for c in range(0, C):

  338.                     for y in range(in_y, in_y + filter_H):

  339.                         for x in range(in_x, in_x + filter_W):

  340.                             if (x < 0 or x >= W or y < 0 or y >= H):

  341.                                 Y[batch_index, row_idx, col_index] = 0

  342.                             else:

  343.                                 Y[batch_index, row_idx,

  344.                                     col_index] = X[batch_index, c, y, x]

  345.                             row_idx += 1

  346.         return Y

  347. 

  348.     def np_conv2d(X, Filter, padding=0, stride=1):

  349.         """Implement a conv2d as a matrix multiply after im2col."""

  350.         filter_outChannel, filter_inChannel, filter_H, filter_W = Filter.shape

  351.         N, C, H, W = X.shape

  352.         assert (H + 2 * padding - filter_H) % stride == 0

  353.         assert (W + 2 * padding - filter_W) % stride == 0

  354.         out_H = (H + 2 * padding - filter_H) // stride + 1

  355.         out_W = (W + 2 * padding - filter_W) // stride + 1

  356. 

  357.         im2col_matrix = im2col(X, filter_H, filter_W, padding, stride)

  358.         filter_matrix = Filter.reshape(filter_outChannel, -1)

  359.         print("shape", im2col_matrix.shape)

  360.         print("shape", filter_matrix.shape)

  361.         print("shape", np.matmul(filter_matrix, im2col_matrix).shape)

  362.         return np.matmul(filter_matrix, im2col_matrix).reshape(N, filter_outChannel, out_H, out_W)

  363.         # return im2col_matrix

  364. 

  365.     shapeX = (100, 3, 28, 28)

  366.     shapeF = (10, 3, 5, 5)

  367.     shapeY = (100, 10, 24, 24)

  368.     shapeW = (100, 3*5*5, 24*24)

  369.     x = np.random.uniform(0, 10, size=shapeX).astype(np.float32)

  370.     f = np.random.uniform(0, 10, size=shapeF).astype(np.float32)

  371.     y = np.zeros(shapeY).astype(np.float32)

  372.     arr_x = ht.array(x, ctx=ctx)

  373.     arr_f = ht.array(f, ctx=ctx)

  374.     arr_y = ht.empty(shapeY, ctx=ctx)

  375.     arr_workspace = ht.empty(shapeW, ctx=ctx)

  376. 

  377.     gpu_op.conv2d(arr_x, arr_f, arr_y, arr_workspace)

  378.     y = arr_y.asnumpy()

  379.     np.testing.assert_allclose(np_conv2d(x, f), y, rtol=1e-5)

  380. 

  381. 

  382. def test_conv2d_Gradient():

  383.     ctx = ht.gpu(0)

  384. 

  385.     def im2col(X, filter_H, filter_W, padding, stride):

  386.         N, C, H, W = X.shape

  387.         assert (H + 2 * padding - filter_H) % stride == 0

  388.         assert (W + 2 * padding - filter_W) % stride == 0

  389.         out_H = (H + 2 * padding - filter_H) // stride + 1

  390.         out_W = (W + 2 * padding - filter_W) // stride + 1

  391. 

  392.         y_row_size = C * filter_H * filter_W

  393.         y_col_size = out_H * out_W

  394.         y_shape = (N, y_row_size, y_col_size)

  395.         Y = np.empty(y_shape, dtype=X.dtype)

  396. 

  397.         for batch_index in range(N):

  398.             for col_index in range(y_col_size):

  399.                 out_y = col_index // out_W

  400.                 out_x = col_index % out_W

  401.                 in_y = out_y * stride - padding

  402.                 in_x = out_x * stride - padding

  403.                 row_idx = 0

  404.                 for c in range(0, C):

  405.                     for y in range(in_y, in_y + filter_H):

  406.                         for x in range(in_x, in_x + filter_W):

  407.                             if (x < 0 or x >= W or y < 0 or y >= H):

  408.                                 Y[batch_index, row_idx, col_index] = 0

  409.                             else:

  410.                                 Y[batch_index, row_idx,

  411.                                     col_index] = X[batch_index, c, y, x]

  412.                             row_idx += 1

  413.         return Y

  414. 

  415.     def np_conv2d(X, Filter, padding=0, stride=1):

  416.         """Implement a conv2d as a matrix multiply after im2col."""

  417.         filter_outChannel, filter_inChannel, filter_H, filter_W = Filter.shape

  418.         N, C, H, W = X.shape

  419.         assert (H + 2 * padding - filter_H) % stride == 0

  420.         assert (W + 2 * padding - filter_W) % stride == 0

  421.         out_H = (H + 2 * padding - filter_H) // stride + 1

  422.         out_W = (W + 2 * padding - filter_W) // stride + 1

  423. 

  424.         im2col_matrix = im2col(X, filter_H, filter_W, padding, stride)

  425.         filter_matrix = Filter.reshape(filter_outChannel, -1)

  426.         return np.matmul(filter_matrix, im2col_matrix).reshape(N, filter_outChannel, out_H, out_W)

  427. 

  428.     def im2col_transpose(X, filter_H, filter_W, Y, padding, stride):

  429.         N, C, H, W = X.shape

  430.         assert (H + 2 * padding - filter_H) % stride == 0

  431.         assert (W + 2 * padding - filter_W) % stride == 0

  432.         out_H = (H + 2 * padding - filter_H) // stride + 1

  433.         out_W = (W + 2 * padding - filter_W) // stride + 1

  434.         _, y_row_size, y_col_size = Y.shape

  435. 

  436.         der_X_shape = (N, C, H, W)

  437.         der_X = np.zeros(der_X_shape, dtype=X.dtype)

  438. 

  439.         for batch_index in range(N):

  440.             for col_index in range(y_col_size):

  441.                 out_y = col_index // out_W

  442.                 out_x = col_index % out_W

  443.                 in_y = out_y * stride - padding

  444.                 in_x = out_x * stride - padding

  445.                 row_idx = 0

  446.                 for c in range(0, C):

  447.                     for y in range(in_y, in_y + filter_H):

  448.                         for x in range(in_x, in_x + filter_W):

  449.                             if (x < 0 or x >= W or y < 0 or y >= H):

  450.                                 Y[batch_index, row_idx, col_index] = 0

  451.                             else:

  452.                                 der_X[batch_index, c, y,

  453.                                       x] += Y[batch_index, row_idx, col_index]

  454.                             row_idx += 1

  455.         return der_X

  456. 

  457.     def np_conv2d_transpose(X, Filter, Y, padding=0, stride=1):

  458.         """Implement a conv2d_transpose as a matrix multiply after im2col."""

  459.         filter_outChannel, filter_inChannel, filter_H, filter_W = Filter.shape

  460.         X_N, X_C, X_H, X_W = X.shape

  461.         Y_N, Y_C, Y_H, Y_W = Y.shape

  462.         YY = Y.reshape((Y_N, Y_C, Y_H * Y_W))    # transformed to im2col Y

  463.         # XX = X.reshape((X_N, X_C, X_W * X_H))   # transformed to im2col X

  464.         F_filter = Filter.reshape((filter_outChannel, -1))

  465.         gradient_im2col_XX = np.matmul(F_filter.T, YY)

  466. 

  467.         gradient_X = im2col_transpose(

  468.             X, filter_H, filter_W, gradient_im2col_XX, padding, stride)    # gradient of x

  469.         im2col_XX = im2col(X, filter_H, filter_W, padding, stride)

  470.         gradient_filter = np.zeros(shape=F_filter.shape, dtype=X.dtype)

  471. 

  472.         for i in range(X_N):

  473.             gradient_filter += np.matmul(YY[i], im2col_XX[i].T)

  474.         gradient_filter = gradient_filter.reshape(Filter.shape)

  475. 

  476.         return gradient_X, gradient_filter

  477. 

  478.     shapeX = (100, 3, 28, 28)

  479.     shapeF = (10, 3, 5, 5)

  480.     shapeY = (100, 10, 24, 24)

  481.     shapeW = (100, 3*5*5, 24*24)

  482.     shapeFF = (100, 10, 3, 5, 5)

  483.     #  input : x , filter : f , output: y

  484.     x = np.random.uniform(0, 10, size=shapeX).astype(np.float32)

  485.     f = np.random.uniform(0, 10, size=shapeF).astype(np.float32)

  486. 

  487.     der_y = np.ones(shape=shapeY)

  488.     gradient_x, gradient_f = np_conv2d_transpose(x, f, der_y)

  489. 

  490.     arr_x = ht.array(x, ctx=ctx)

  491.     arr_f = ht.array(f, ctx=ctx)

  492.     gradient_y = ht.array(der_y, ctx=ctx)

  493.     gradient_xx = ht.array(x, ctx=ctx)

  494.     gradient_ff = ht.array(f, ctx=ctx)

  495. 

  496.     arr_workspace_im2col = ht.empty(shapeW, ctx=ctx)

  497.     arr_workspace_batch_filter = ht.empty(shapeFF, ctx=ctx)

  498.     gpu_op.conv2d_gradient_of_filter(

  499.         arr_x, gradient_y, gradient_ff, arr_workspace_im2col, arr_workspace_batch_filter)

  500.     gpu_op.conv2d_gradient_of_data(

  501.         arr_f, gradient_y, gradient_xx, arr_workspace_im2col)

  502. 

  503.     np.testing.assert_allclose(gradient_x, gradient_xx.asnumpy(), rtol=1e-5)

  504.     # test ok

  505.     np.testing.assert_allclose(gradient_f, gradient_ff.asnumpy(), rtol=1e-5)

  506. 

  507. 

  508. def test_average_pooling():

  509.     ctx = ht.gpu(0)

  510. 

  511.     def np_average_pooling(input, kernel_H, kernel_W, padding=0, stride=1):

  512.         N, C, H, W = input.shape

  513.         assert((H + 2 * padding - kernel_H) % stride == 0)

  514.         assert((W + 2 * padding - kernel_W) % stride == 0)

  515.         pooled_H = (H + 2 * padding - kernel_H) // stride + 1

  516.         pooled_W = (W + 2 * padding - kernel_W) // stride + 1

  517.         pooled_layer = np.zeros(

  518.             shape=(N, C, pooled_H, pooled_W), dtype=np.float32)

  519.         pooling_size = kernel_H * kernel_W

  520.         for n in range(N):

  521.             for c in range(C):

  522.                 for h in range(pooled_H):

  523.                     for w in range(pooled_W):

  524.                         hs = h * stride - padding

  525.                         ws = w * stride - padding

  526.                         hend = min(hs + kernel_H, H)

  527.                         wend = min(ws + kernel_W, W)

  528.                         hs = max(hs, 0)

  529.                         ws = max(ws, 0)

  530.                         for i in range(hs, hend):

  531.                             for j in range(ws, wend):

  532.                                 pooled_layer[n][c][h][w] += input[n][c][i][j]

  533.                         pooled_layer[n][c][h][w] /= pooling_size

  534.         return pooled_layer

  535. 

  536.     def np_average_pooling_gradient(gradient_y, kernel_H, kernel_W, padding=0, stride=1):

  537.         N, C, pooled_H, pooled_W = gradient_y.shape

  538.         H = (pooled_H - 1) * stride + kernel_H - 2 * padding

  539.         W = (pooled_W - 1) * stride + kernel_W - 2 * padding

  540. 

  541.         gradient_x = np.zeros(shape=(N, C, H, W), dtype=np.float32)

  542.         pooling_size = kernel_H * kernel_W

  543.         for n in range(N):

  544.             for c in range(C):

  545.                 for h in range(pooled_H):

  546.                     for w in range(pooled_W):

  547.                         hs = h * stride - padding

  548.                         ws = w * stride - padding

  549.                         hend = min(hs + kernel_H, H)

  550.                         wend = min(ws + kernel_W, W)

  551.                         hs = max(hs, 0)

  552.                         ws = max(ws, 0)

  553.                         for i in range(hs, hend):

  554.                             for j in range(ws, wend):

  555.                                 gradient_x[n][c][i][j] += gradient_y[n][c][h][w] / \

  556.                                     pooling_size

  557. 

  558.         return gradient_x

  559. 

  560.     shapeX = (100, 3, 28, 28)

  561.     # (1,1,5,5)

  562.     shapeY = (100, 3, 24, 24)

  563.     #  input : x , filter : f , output: y

  564.     x = np.random.uniform(0, 10, size=shapeX).astype(np.float32)

  565.     gradient_y = np.random.uniform(0, 10, size=shapeY).astype(np.float32)

  566. 

  567.     arr_x = ht.array(x, ctx=ctx)

  568.     arr_gradient_y = ht.array(gradient_y, ctx=ctx)

  569.     arr_pool_layer = ht.empty(shapeY, ctx=ctx)

  570.     arr_gradient_x = ht.empty(shapeX, ctx=ctx)

  571. 

  572.     gpu_op.average_pooling2d(arr_x, 5, 5, arr_pool_layer)

  573.     gpu_op.average_pooling2d_gradient(arr_gradient_y, 5, 5, arr_gradient_x)

  574. 

  575.     np_pool_layer = np_average_pooling(x, 5, 5)

  576.     np_gradient_x = np_average_pooling_gradient(gradient_y, 5, 5)

  577. 

  578.     np.testing.assert_allclose(

  579.         np_pool_layer, arr_pool_layer.asnumpy(), rtol=1e-5)

  580. 

  581.     np.testing.assert_allclose(

  582.         np_gradient_x, arr_gradient_x.asnumpy(), rtol=1e-5)

  583. 

  584. 

  585. def test_reshape():

  586.     ctx = ht.gpu(0)

  587. 

  588.     def np_reshape(X, output_shape):

  589.         return X.reshape(output_shape)

  590. 

  591.     shapeX = (10, 5, 28, 28)

  592.     shapeY = (50, 28, 28)

  593. 

  594.     x = np.random.uniform(0, 10, size=shapeX).astype(np.float32)

  595.     y = np_reshape(x, shapeY)

  596. 

  597.     arr_x = ht.array(x, ctx=ctx)

  598.     arr_y = ht.empty(shapeY, ctx=ctx)

  599.     gpu_op.array_reshape(arr_x, arr_y)

  600.     np.testing.assert_allclose(y, arr_y.asnumpy(), rtol=1e-5)

  601. 

  602. 

  603. def test_conv2d_broadcast_to():

  604.     ctx = ht.gpu(0)

  605.     shapeX = (32)

  606.     shapeY = (100, 32, 28, 28)

  607.     shapeW = (100, 28, 28, 32)

  608.     x = np.random.uniform(0, 10, size=shapeX).astype(np.float32)

  609.     np_y = np.broadcast_to(x, shapeW)

  610.     np_y = np_y.swapaxes(1, 3)

  611. 

  612.     arr_x = ht.array(x, ctx=ctx)

  613.     arr_y = ht.empty(shapeY, ctx=ctx)

  614.     gpu_op.conv2d_broadcast_to(arr_x, arr_y)

  615. 

  616.     np.testing.assert_allclose(np_y, arr_y.asnumpy(), rtol=1e-5)

  617. 

  618. 

  619. def test_conv2d_reduce_sum():

  620.     ctx = ht.gpu(0)

  621.     shapeX = (32,)

  622.     shapeY = (100, 32, 28, 28)

  623.     shapeW = (100, 28, 28, 32)

  624.     x = np.random.uniform(0, 10, size=shapeY).astype(np.float32)

  625.     np_y = np.sum(x, axis=(0, 2, 3))

  626. 

  627.     arr_x = ht.array(x, ctx=ctx)

  628.     arr_y = ht.empty(shapeX, ctx=ctx)

  629. 

  630.     gpu_op.conv2d_reduce_sum(arr_x, arr_y)

  631.     np.testing.assert_allclose(np_y, arr_y.asnumpy(), rtol=1e-5)

  632. 

  633. 

  634. def test_cudnn_conv2d():

  635.     ctx = ht.gpu(0)

  636.     # im2col and np_conv2d are helper functions

  637. 

  638.     def im2col(X, filter_H, filter_W, padding, stride):

  639.         N, C, H, W = X.shape

  640.         assert (H + 2 * padding - filter_H) % stride == 0

  641.         assert (W + 2 * padding - filter_W) % stride == 0

  642.         out_H = (H + 2 * padding - filter_H) // stride + 1

  643.         out_W = (W + 2 * padding - filter_W) // stride + 1

  644. 

  645.         y_row_size = C * filter_H * filter_W

  646.         y_col_size = out_H * out_W

  647.         y_shape = (N, y_row_size, y_col_size)

  648.         Y = np.empty(y_shape, dtype=X.dtype)

  649. 

  650.         for batch_index in range(N):

  651.             for col_index in range(y_col_size):

  652.                 out_y = col_index // out_W

  653.                 out_x = col_index % out_W

  654.                 in_y = out_y * stride - padding

  655.                 in_x = out_x * stride - padding

  656.                 row_idx = 0

  657.                 for c in range(0, C):

  658.                     for y in range(in_y, in_y + filter_H):

  659.                         for x in range(in_x, in_x + filter_W):

  660.                             if (x < 0 or x >= W or y < 0 or y >= H):

  661.                                 Y[batch_index, row_idx, col_index] = 0

  662.                             else:

  663.                                 Y[batch_index, row_idx,

  664.                                     col_index] = X[batch_index, c, y, x]

  665.                             row_idx += 1

  666.         return Y

  667. 

  668.     def np_conv2d(X, Filter, padding=0, stride=1):

  669.         """Implement a conv2d as a matrix multiply after im2col."""

  670.         filter_outChannel, filter_inChannel, filter_H, filter_W = Filter.shape

  671.         N, C, H, W = X.shape

  672.         assert (H + 2 * padding - filter_H) % stride == 0

  673.         assert (W + 2 * padding - filter_W) % stride == 0

  674.         out_H = (H + 2 * padding - filter_H) // stride + 1

  675.         out_W = (W + 2 * padding - filter_W) // stride + 1

  676. 

  677.         im2col_matrix = im2col(X, filter_H, filter_W, padding, stride)

  678.         filter_matrix = Filter.reshape(filter_outChannel, -1)

  679.         print("shape", im2col_matrix.shape)

  680.         print("shape", filter_matrix.shape)

  681.         print("shape", np.matmul(filter_matrix, im2col_matrix).shape)

  682.         return np.matmul(filter_matrix, im2col_matrix).reshape(N, filter_outChannel, out_H, out_W)

  683.         # return im2col_matrix

  684. 

  685.     shapeX = (100, 3, 28, 28)

  686.     shapeF = (10, 3, 5, 5)

  687.     shapeY = (100, 10, 24, 24)

  688.     shapeW = (100, 3*5*5, 24*24)

  689.     x = np.random.uniform(0, 10, size=shapeX).astype(np.float32)

  690.     f = np.random.uniform(0, 10, size=shapeF).astype(np.float32)

  691.     y = np.zeros(shapeY).astype(np.float32)

  692.     arr_x = ht.array(x, ctx=ctx)

  693.     arr_f = ht.array(f, ctx=ctx)

  694.     arr_y = ht.empty(shapeY, ctx=ctx)

  695.     arr_workspace = ht.empty(shapeW, ctx=ctx)

  696. 

  697.     gpu_op.CuDNN_conv2d(arr_x, arr_f, arr_y)

  698.     y = arr_y.asnumpy()

  699.     np.testing.assert_allclose(np_conv2d(x, f), y, rtol=1e-5)

  700. 

  701. 

  702. def test_cudnn_conv2d_Gradient():

  703.     ctx = ht.gpu(0)

  704. 

  705.     def im2col(X, filter_H, filter_W, padding, stride):

  706.         N, C, H, W = X.shape

  707.         assert (H + 2 * padding - filter_H) % stride == 0

  708.         assert (W + 2 * padding - filter_W) % stride == 0

  709.         out_H = (H + 2 * padding - filter_H) // stride + 1

  710.         out_W = (W + 2 * padding - filter_W) // stride + 1

  711. 

  712.         y_row_size = C * filter_H * filter_W

  713.         y_col_size = out_H * out_W

  714.         y_shape = (N, y_row_size, y_col_size)

  715.         Y = np.empty(y_shape, dtype=X.dtype)

  716. 

  717.         for batch_index in range(N):

  718.             for col_index in range(y_col_size):

  719.                 out_y = col_index // out_W

  720.                 out_x = col_index % out_W

  721.                 in_y = out_y * stride - padding

  722.                 in_x = out_x * stride - padding

  723.                 row_idx = 0

  724.                 for c in range(0, C):

  725.                     for y in range(in_y, in_y + filter_H):

  726.                         for x in range(in_x, in_x + filter_W):

  727.                             if (x < 0 or x >= W or y < 0 or y >= H):

  728.                                 Y[batch_index, row_idx, col_index] = 0

  729.                             else:

  730.                                 Y[batch_index, row_idx,

  731.                                     col_index] = X[batch_index, c, y, x]

  732.                             row_idx += 1

  733.         return Y

  734. 

  735.     def np_conv2d(X, Filter, padding=0, stride=1):

  736.         """Implement a conv2d as a matrix multiply after im2col."""

  737.         filter_outChannel, filter_inChannel, filter_H, filter_W = Filter.shape

  738.         N, C, H, W = X.shape

  739.         assert (H + 2 * padding - filter_H) % stride == 0

  740.         assert (W + 2 * padding - filter_W) % stride == 0

  741.         out_H = (H + 2 * padding - filter_H) // stride + 1

  742.         out_W = (W + 2 * padding - filter_W) // stride + 1

  743. 

  744.         im2col_matrix = im2col(X, filter_H, filter_W, padding, stride)

  745.         filter_matrix = Filter.reshape(filter_outChannel, -1)

  746.         return np.matmul(filter_matrix, im2col_matrix).reshape(N, filter_outChannel, out_H, out_W)

  747. 

  748.     def im2col_transpose(X, filter_H, filter_W, Y, padding, stride):

  749.         N, C, H, W = X.shape

  750.         assert (H + 2 * padding - filter_H) % stride == 0

  751.         assert (W + 2 * padding - filter_W) % stride == 0

  752.         out_H = (H + 2 * padding - filter_H) // stride + 1

  753.         out_W = (W + 2 * padding - filter_W) // stride + 1

  754.         _, y_row_size, y_col_size = Y.shape

  755. 

  756.         der_X_shape = (N, C, H, W)

  757.         der_X = np.zeros(der_X_shape, dtype=X.dtype)

  758. 

  759.         for batch_index in range(N):

  760.             for col_index in range(y_col_size):

  761.                 out_y = col_index // out_W

  762.                 out_x = col_index % out_W

  763.                 in_y = out_y * stride - padding

  764.                 in_x = out_x * stride - padding

  765.                 row_idx = 0

  766.                 for c in range(0, C):

  767.                     for y in range(in_y, in_y + filter_H):

  768.                         for x in range(in_x, in_x + filter_W):

  769.                             if (x < 0 or x >= W or y < 0 or y >= H):

  770.                                 Y[batch_index, row_idx, col_index] = 0

  771.                             else:

  772.                                 der_X[batch_index, c, y,

  773.                                       x] += Y[batch_index, row_idx, col_index]

  774.                             row_idx += 1

  775.         return der_X

  776. 

  777.     def np_conv2d_transpose(X, Filter, Y, padding=0, stride=1):

  778.         """Implement a conv2d_transpose as a matrix multiply after im2col."""

  779.         filter_outChannel, filter_inChannel, filter_H, filter_W = Filter.shape

  780.         X_N, X_C, X_H, X_W = X.shape

  781.         Y_N, Y_C, Y_H, Y_W = Y.shape

  782.         YY = Y.reshape((Y_N, Y_C, Y_H * Y_W))    # transformed to im2col Y

  783.         # XX = X.reshape((X_N, X_C, X_W * X_H))   # transformed to im2col X

  784.         F_filter = Filter.reshape((filter_outChannel, -1))

  785.         gradient_im2col_XX = np.matmul(F_filter.T, YY)

  786. 

  787.         gradient_X = im2col_transpose(

  788.             X, filter_H, filter_W, gradient_im2col_XX, padding, stride)    # gradient of x

  789.         im2col_XX = im2col(X, filter_H, filter_W, padding, stride)

  790.         gradient_filter = np.zeros(shape=F_filter.shape, dtype=X.dtype)

  791. 

  792.         for i in range(X_N):

  793.             gradient_filter += np.matmul(YY[i], im2col_XX[i].T)

  794.         gradient_filter = gradient_filter.reshape(Filter.shape)

  795. 

  796.         return gradient_X, gradient_filter

  797. 

  798.     shapeX = (100, 3, 28, 28)

  799.     shapeF = (10, 3, 5, 5)

  800.     shapeY = (100, 10, 24, 24)

  801.     shapeW = (100, 3*5*5, 24*24)

  802.     shapeFF = (100, 10, 3, 5, 5)

  803.     #  input : x , filter : f , output: y

  804.     x = np.random.uniform(0, 10, size=shapeX).astype(np.float32)

  805.     f = np.random.uniform(0, 10, size=shapeF).astype(np.float32)

  806. 

  807.     der_y = np.ones(shape=shapeY)

  808.     gradient_x, gradient_f = np_conv2d_transpose(x, f, der_y)

  809. 

  810.     arr_x = ht.array(x, ctx=ctx)

  811.     arr_f = ht.array(f, ctx=ctx)

  812.     gradient_y = ht.array(der_y, ctx=ctx)

  813.     gradient_xx = ht.array(x, ctx=ctx)

  814.     gradient_ff = ht.array(f, ctx=ctx)

  815. 

  816.     arr_workspace_im2col = ht.empty(shapeW, ctx=ctx)

  817.     arr_workspace_batch_filter = ht.empty(shapeFF, ctx=ctx)

  818.     gpu_op.CuDNN_conv2d_gradient_of_filter(arr_x, gradient_y, gradient_ff)

  819.     gpu_op.CuDNN_conv2d_gradient_of_data(arr_f, gradient_y, gradient_xx)

  820. 

  821.     np.testing.assert_allclose(gradient_x, gradient_xx.asnumpy(), rtol=1e-5)

  822.     # test ok

  823.     np.testing.assert_allclose(gradient_f, gradient_ff.asnumpy(), rtol=1e-5)

  824. 

  825. 

  826. def test_average_pooling():

  827.     ctx = ht.gpu(0)

  828. 

  829.     def np_average_pooling(input, kernel_H, kernel_W, padding=0, stride=1):

  830.         N, C, H, W = input.shape

  831.         assert((H + 2 * padding - kernel_H) % stride == 0)

  832.         assert((W + 2 * padding - kernel_W) % stride == 0)

  833.         pooled_H = (H + 2 * padding - kernel_H) // stride + 1

  834.         pooled_W = (W + 2 * padding - kernel_W) // stride + 1

  835.         pooled_layer = np.zeros(

  836.             shape=(N, C, pooled_H, pooled_W), dtype=np.float32)

  837.         pooling_size = kernel_H * kernel_W

  838.         for n in range(N):

  839.             for c in range(C):

  840.                 for h in range(pooled_H):

  841.                     for w in range(pooled_W):

  842.                         hs = h * stride - padding

  843.                         ws = w * stride - padding

  844.                         hend = min(hs + kernel_H, H)

  845.                         wend = min(ws + kernel_W, W)

  846.                         hs = max(hs, 0)

  847.                         ws = max(ws, 0)

  848.                         for i in range(hs, hend):

  849.                             for j in range(ws, wend):

  850.                                 pooled_layer[n][c][h][w] += input[n][c][i][j]

  851.                         pooled_layer[n][c][h][w] /= pooling_size

  852.         return pooled_layer

  853. 

  854.     def np_average_pooling_gradient(gradient_y, kernel_H, kernel_W, padding=0, stride=1):

  855.         N, C, pooled_H, pooled_W = gradient_y.shape

  856.         H = (pooled_H - 1) * stride + kernel_H - 2 * padding

  857.         W = (pooled_W - 1) * stride + kernel_W - 2 * padding

  858. 

  859.         gradient_x = np.zeros(shape=(N, C, H, W), dtype=np.float32)

  860.         pooling_size = kernel_H * kernel_W

  861.         for n in range(N):

  862.             for c in range(C):

  863.                 for h in range(pooled_H):

  864.                     for w in range(pooled_W):

  865.                         hs = h * stride - padding

  866.                         ws = w * stride - padding

  867.                         hend = min(hs + kernel_H, H)

  868.                         wend = min(ws + kernel_W, W)

  869.                         hs = max(hs, 0)

  870.                         ws = max(ws, 0)

  871.                         for i in range(hs, hend):

  872.                             for j in range(ws, wend):

  873.                                 gradient_x[n][c][i][j] += gradient_y[n][c][h][w] / \

  874.                                     pooling_size

  875. 

  876.         return gradient_x

  877. 

  878.     shapeX = (100, 3, 28, 28)

  879.     shapeY = (100, 3, 24, 24)

  880.     x = np.random.uniform(0, 10, size=shapeX).astype(np.float32)

  881.     gradient_y = np.random.uniform(0, 10, size=shapeY).astype(np.float32)

  882.     arr_x = ht.array(x, ctx=ctx)

  883.     arr_gradient_y = ht.array(gradient_y, ctx=ctx)

  884.     arr_pool_layer = ht.empty(shapeY, ctx=ctx)

  885.     arr_gradient_x = ht.empty(shapeX, ctx=ctx)

  886. 

  887.     gpu_op.CuDNN_average_pooling2d(arr_x, 5, 5, arr_pool_layer)

  888.     gpu_op.CuDNN_average_pooling2d_gradient(

  889.         arr_pool_layer, arr_gradient_y, arr_x, 5, 5, arr_gradient_x)

  890.     np_pool_layer = np_average_pooling(x, 5, 5)

  891.     np_gradient_x = np_average_pooling_gradient(gradient_y, 5, 5)

  892. 

  893.     np.testing.assert_allclose(

  894.         np_pool_layer, arr_pool_layer.asnumpy(), rtol=1e-5)

  895. 

  896.     np.testing.assert_allclose(

  897.         np_gradient_x, arr_gradient_x.asnumpy(), rtol=1e-5)

  898. 

  899. 

  900. def test_CuDNN_max_pooling():

  901.     ctx = ht.gpu(0)

  902. 

  903.     def np_max_pooling_gradient(input, gradient_y, kernel_H, kernel_W, padding=0, stride=1):

  904.         N, C, pooled_H, pooled_W = gradient_y.shape

  905.         H = (pooled_H - 1) * stride + kernel_H - 2 * padding

  906.         W = (pooled_W - 1) * stride + kernel_W - 2 * padding

  907.         gradient_x = np.zeros(shape=(N, C, H, W), dtype=np.float32)

  908.         pooling_size = kernel_H * kernel_W

  909. 

  910.         for n in range(N):

  911.             for c in range(C):

  912.                 for h in range(pooled_H):

  913.                     for w in range(pooled_W):

  914.                         hs = h * stride - padding

  915.                         ws = w * stride - padding

  916.                         hend = min(hs + kernel_H, H)

  917.                         wend = min(ws + kernel_W, W)

  918.                         hs = max(hs, 0)

  919.                         ws = max(ws, 0)

  920. 

  921.                         hargmax = hs

  922.                         wargmax = ws

  923.                         for i in range(hs, hend):

  924.                             for j in range(ws, wend):

  925.                                 if input[n][c][i][j] > input[n][c][hargmax][wargmax]:

  926.                                     hargmax = i

  927.                                     wargmax = j

  928.                         gradient_x[n][c][hargmax][wargmax] += gradient_y[n][c][h][w]

  929. 

  930.         return gradient_x

  931. 

  932.     shapeX = (100, 3, 28, 28)

  933.     # # (1,1,5,5)

  934.     shapeY = (100, 3, 24, 24)

  935.     x = np.random.uniform(0, 10, size=shapeX).astype(np.float32)

  936.     gradient_y = np.random.uniform(0, 10, size=shapeY).astype(np.float32)

  937.     arr_x = ht.array(x, ctx=ctx)

  938.     arr_gradient_y = ht.array(gradient_y, ctx=ctx)

  939.     arr_pool_layer = ht.empty(shapeY, ctx=ctx)

  940.     arr_gradient_x = ht.empty(shapeX, ctx=ctx)

  941. 

  942.     arr_pool_layer1 = ht.empty(shapeY, ctx=ctx)

  943.     arr_gradient_x1 = ht.empty(shapeX, ctx=ctx)

  944. 

  945.     gpu_op.CuDNN_max_pooling2d(arr_x, 2, 2, arr_pool_layer)

  946.     gpu_op.CuDNN_max_pooling2d_gradient(

  947.         arr_pool_layer, arr_gradient_y, arr_x, 2, 2, arr_gradient_x)

  948. 

  949.     gpu_op.max_pooling2d(arr_x, 2, 2, arr_pool_layer1)

  950.     gpu_op.max_pooling2d_gradient(arr_x, arr_gradient_y, 2, 2, arr_gradient_x1)

  951. 

  952.     np.testing.assert_allclose(

  953.         arr_pool_layer.asnumpy(), arr_pool_layer1.asnumpy(), rtol=1e-5)

  954. 

  955.     np.testing.assert_allclose(

  956.         arr_gradient_x.asnumpy(), arr_gradient_x1.asnumpy(), rtol=1e-5)

  957. 

  958. 

  959. def test_CuDNN_dropout_op():

  960.     import ctypes

  961.     ctx = ht.gpu(0)

  962.     shapeX = (3, 10)

  963.     x = np.random.uniform(0, 10, size=shapeX).astype(np.float32)

  964.     arr_x = ht.array(x, ctx=ctx)

  965.     arr_y = ht.empty(shapeX, ctx=ctx)

  966.     shapeK = (1)

  967.     keep_prob = 0.5

  968.     reserve_size = ctypes.c_int(0)

  969.     reserve_space = ctypes.c_void_p(0)

  970.     gpu_op.CuDNN_Dropout(arr_x, keep_prob, arr_y,

  971.                          reserve_size, reserve_space, 1)

  972. 

  973.     gradient_y = np.random.uniform(0, 10, size=shapeX).astype(np.float32)

  974.     gradient_y = ht.array(gradient_y, ctx=ctx)

  975.     gradient_x = ht.empty(shapeX, ctx=ctx)

  976.     gpu_op.CuDNN_Dropout_gradient(

  977.         gradient_y, keep_prob, gradient_x, reserve_size, reserve_space)

  978.     print(arr_y.asnumpy())

  979.     print(gradient_x.asnumpy())

  980. 

  981. 

  982. def test_pad():

  983.     ctx = ht.gpu(0)

  984.     shape = (1, 1, 1, 3)

  985.     paddings = [[1, 1], [1, 1]]

  986.     to_shape = (1, 1, 3, 5)

  987.     x = np.random.uniform(0, 10, size=shape).astype(np.float32)

  988.     arr_x = ht.array(x, ctx=ctx)

  989.     arr_y = ht.empty(to_shape, ctx=ctx)

  990.     gpu_op.pad(arr_x, arr_y, paddings)

  991.     print(arr_x.asnumpy())

  992.     print(arr_y.asnumpy())

  993.     gradient_y = np.random.uniform(0, 10, size=to_shape).astype(np.float32)

  994.     arr_gradient_y = ht.array(gradient_y, ctx=ctx)

  995.     arr_gradient_x = ht.empty(shape, ctx=ctx)

  996.     gpu_op.pad_gradient(arr_gradient_y, arr_gradient_x, paddings)

  997.     print(arr_gradient_y.asnumpy())

  998.     print(arr_gradient_x.asnumpy())

  999. 

  1000. 

  1001. def test_concat():

  1002.     def unit_test(shape1, shape2, axis):

  1003.         ctx = ht.gpu(0)

  1004.         x1 = np.random.random(shape1).astype(np.float32)

  1005.         x2 = np.random.random(shape2).astype(np.float32)

  1006.         arr_x1 = ht.array(x1, ctx=ctx)

  1007.         arr_x2 = ht.array(x2, ctx=ctx)

  1008.         np_res = np.concatenate([x1, x2], axis)

  1009.         arr_res = ht.empty(np_res.shape, ctx=ctx)

  1010.         gpu_op.concat(arr_x1, arr_x2, arr_res, axis)

  1011.         np.testing.assert_allclose(arr_res.asnumpy(), np_res)

  1012. 

  1013.         grad_x1 = ht.empty(shape1, ctx=ctx)

  1014.         grad_x2 = ht.empty(shape2, ctx=ctx)

  1015.         gpu_op.concat_gradient(arr_res, grad_x1, axis=axis, idx=0)

  1016.         gpu_op.concat_gradient(arr_res, grad_x2, axis=axis, idx=1)

  1017.         np.testing.assert_allclose(x1, grad_x1.asnumpy())

  1018.         np.testing.assert_allclose(x2, grad_x2.asnumpy())

  1019. 

  1020.         print("Pass test with ", shape1, shape2, axis)

  1021. 

  1022.     unit_test((1, 2), (1, 2), 0)

  1023.     unit_test((12, 34, 56), (12, 43, 56), 1)

  1024. 

  1025. 

  1026. def test_matrix_transpose():

  1027.     shape = (4321, 1234)

  1028.     ctx = ht.gpu(0)

  1029.     x = np.random.uniform(-1, 1, shape).astype(np.float32)

  1030.     arr_x = ht.array(x, ctx=ctx)

  1031.     arr_y = ht.empty((shape[1], shape[0]), ctx=ctx)

  1032.     gpu_op.matrix_transpose(arr_x, arr_y, perm=[1, 0])

  1033.     y = arr_y.asnumpy()

  1034.     np.testing.assert_allclose(np.transpose(x), y)

  1035. 

  1036.     shape = (21, 43, 65, 11)

  1037.     x = np.random.uniform(-1, 1, shape).astype(np.float32)

  1038.     arr_x = ht.array(x, ctx=ctx)

  1039.     arr_y = ht.empty((65, 11, 43, 21), ctx=ctx)

  1040.     gpu_op.matrix_transpose(arr_x, arr_y, perm=[2, 3, 1, 0])

  1041.     y = arr_y.asnumpy()

  1042.     np.testing.assert_allclose(np.transpose(x, [2, 3, 1, 0]), y)

  1043. 

  1044. 

  1045. def test_slice():

  1046.     i_shape = (123, 234, 13, 7)

  1047.     o_shape = (67, 209, 3, 5)

  1048.     begin_pos = (31, 11, 3, 1)

  1049. 

  1050.     ctx = ht.gpu(0)

  1051.     x = np.random.normal(size=i_shape).astype(np.float32)

  1052.     arr_x = ht.array(x, ctx=ctx)

  1053.     arr_y = ht.empty(o_shape, ctx=ctx)

  1054. 

  1055.     gpu_op.matrix_slice(arr_x, arr_y, begin_pos)

  1056.     y = arr_y.asnumpy()

  1057.     index = tuple([slice(i, i+j) for i, j in zip(begin_pos, o_shape)])

  1058.     np.testing.assert_allclose(x[index], y)

  1059.     print("Slice op no bug.")

  1060. 

  1061.     begin_pos = (29, 3, 5, 0)

  1062.     gpu_op.matrix_slice_gradient(arr_y, arr_x, begin_pos)

  1063.     x = arr_x.asnumpy()

  1064.     index = tuple([slice(i, i+j) for i, j in zip(begin_pos, o_shape)])

  1065.     x_ = np.zeros(i_shape, dtype=np.float32)

  1066.     x_[index] = y

  1067.     np.testing.assert_allclose(x, x_)

  1068.     print("Slice gradient op no bug.")

  1069. 

  1070. 

  1071. def test_where():

  1072.     cond = np.random.randint(2, size=(5, 5)).astype(np.float32)

  1073.     x = np.random.rand(5, 5).astype(np.float32)

  1074.     y = np.random.rand(5, 5).astype(np.float32)

  1075.     ctx = ht.gpu(0)

  1076.     cond_gpu = ht.array(cond, ctx=ctx)

  1077.     x_gpu = ht.array(x, ctx=ctx)

  1078.     y_gpu = ht.array(y, ctx=ctx)

  1079.     output_gpu = ht.empty((5, 5), ctx=ctx)

  1080. 

  1081.     gpu_op.where(cond_gpu, x_gpu, y_gpu, output_gpu)

  1082.     output = np.where(cond, x, y)

  1083. 

  1084.     np.testing.assert_allclose(output, output_gpu.asnumpy())

  1085. 

  1086. 

  1087. def test_batch_matrix_multiply():

  1088.     ctx = ht.gpu(0)

  1089.     x = np.random.uniform(0, 10, size=(2, 3, 500, 700)).astype(np.float32)

  1090.     y = np.random.uniform(0, 10, size=(2, 3, 700, 1000)).astype(np.float32)

  1091.     arr_x = ht.array(x, ctx=ctx)

  1092.     arr_y = ht.array(y, ctx=ctx)

  1093.     arr_z = ht.empty((2, 3, 500, 1000), ctx=ctx)

  1094.     gpu_op.batch_matrix_multiply(arr_x, False, arr_y, False, arr_z)

  1095.     z = arr_z.asnumpy()

  1096.     np.testing.assert_allclose(np.matmul(x, y), z, rtol=1e-5)

  1097.     print('Test 1 passed.')

  1098. 

  1099.     x = np.random.uniform(0, 10, size=(7, 11, 1000, 500)).astype(np.float32)

  1100.     y = np.random.uniform(0, 10, size=(7, 11, 2000, 500)).astype(np.float32)

  1101.     arr_x = ht.array(x, ctx=ctx)

  1102.     arr_y = ht.array(y, ctx=ctx)

  1103.     arr_z = ht.empty((7, 11, 1000, 2000), ctx=ctx)

  1104.     gpu_op.batch_matrix_multiply(arr_x, False, arr_y, True, arr_z)

  1105.     z = arr_z.asnumpy()

  1106.     np.testing.assert_allclose(

  1107.         np.matmul(x, np.transpose(y, [0, 1, 3, 2])), z, rtol=1e-5)

  1108.     print('Test 2 passed.')

  1109. 

  1110.     x = np.random.uniform(0, 10, size=(3, 2, 5, 500, 1000)).astype(np.float32)

  1111.     y = np.random.uniform(0, 10, size=(3, 2, 5, 2000, 500)).astype(np.float32)

  1112.     arr_x = ht.array(x, ctx=ctx)

  1113.     arr_y = ht.array(y, ctx=ctx)

  1114.     arr_z = ht.empty((3, 2, 5, 1000, 2000), ctx=ctx)

  1115.     gpu_op.batch_matrix_multiply(arr_x, True, arr_y, True, arr_z)

  1116.     z = arr_z.asnumpy()

  1117.     np.testing.assert_allclose(np.matmul(np.transpose(x, [0, 1, 2, 4, 3]), np.transpose(y, [0, 1, 2, 4, 3])), z,

  1118.                                rtol=1e-5)

  1119.     print('Test 3 passed.')

  1120. 

  1121. 

  1122. def test_broadcast_shape():

  1123.     def unit_test(shape1, shape2):

  1124.         ctx = ht.gpu(0)

  1125.         x = np.random.random(shape1).astype(np.float32)

  1126.         arr_x = ht.array(x, ctx=ctx)

  1127.         arr_y = ht.empty(shape2, ctx=ctx)

  1128.         gpu_op.broadcast_shape(arr_x, arr_y)

  1129.         np.testing.assert_allclose(arr_y.asnumpy(), np.broadcast_to(x, shape2))

  1130.         print('Passed test with input shape %s and output shape %s.' %

  1131.               (str(shape1), str(shape2)))

  1132. 

  1133.     unit_test((3, 1), (2, 3, 4))

  1134.     unit_test((1,), (2, 3, 4, 5))

  1135.     unit_test((1, 1, 3, 1), (9, 8, 3, 7))

  1136. 

  1137. 

  1138. def test_reduce_sum():

  1139.     def unit_test(shape, axes):

  1140.         ctx = ht.gpu(0)

  1141.         x = np.random.random(shape).astype(np.float32)

  1142.         arr_x = ht.array(x, ctx=ctx)

  1143.         o_shape = list(shape)

  1144.         for ax in axes:

  1145.             o_shape[ax] = 0

  1146.         o_shape = [i for i in o_shape if i > 0]

  1147.         arr_y = ht.empty(o_shape, ctx=ctx)

  1148.         gpu_op.reduce_sum(arr_x, arr_y, axes)

  1149.         np.testing.assert_allclose(arr_y.asnumpy(), np.sum(

  1150.             x, tuple(axes), keepdims=False), rtol=1e-6)

  1151.         print('Passed test with input shape %s and reduce axes %s.' %

  1152.               (str(shape), str(axes)))

  1153. 

  1154.     unit_test((2, 3, 4), [2])

  1155.     unit_test((2, 3, 4), [2, 1])

  1156.     unit_test((2, 3, 4), [2, 1, 0])

  1157.     unit_test((2, 3, 1, 5, 6), [1, 2, 4])

  1158. 

  1159. 

  1160. def test_reduce_mean():

  1161.     def unit_test(shape, axes):

  1162.         ctx = ht.gpu(0)

  1163.         x = np.random.random(shape).astype(np.float32)

  1164.         arr_x = ht.array(x, ctx=ctx)

  1165.         o_shape = list(shape)

  1166.         for ax in axes:

  1167.             o_shape[ax] = 0

  1168.         o_shape = [i for i in o_shape if i > 0]

  1169.         arr_y = ht.empty(o_shape, ctx=ctx)

  1170.         gpu_op.reduce_mean(arr_x, arr_y, axes)

  1171.         np.testing.assert_allclose(arr_y.asnumpy(), np.mean(

  1172.             x, tuple(axes), keepdims=False), rtol=1e-6)

  1173.         print('Passed test with input shape %s and reduce axes %s.' %

  1174.               (str(shape), str(axes)))

  1175. 

  1176.     unit_test((2, 3, 4), [2])

  1177.     unit_test((2, 3, 4), [2, 1])

  1178.     unit_test((2, 3, 4), [2, 1, 0])

  1179.     unit_test((2, 3, 1, 5, 6), [1, 2, 4])

  1180. 

  1181. 

  1182. def test_dropout():

  1183.     ctx = ht.gpu(0)

  1184.     # shapeX = (16, 8, 99, 64)

  1185.     shapeX = (10,)

  1186.     x = np.random.uniform(0, 10, size=shapeX).astype(np.float32)

  1187.     arr_x = ht.array(x, ctx=ctx)

  1188.     arr_y = ht.empty(shapeX, ctx=ctx)

  1189.     dropout_rate = 0.6

  1190.     import ctypes

  1191.     seed = ctypes.c_ulonglong(0)

  1192.     print(x)

  1193.     gpu_op.dropout(arr_x, dropout_rate, arr_y, seed)

  1194.     print(arr_y.asnumpy())

  1195.     print(seed)

  1196.     gpu_op.dropout_gradient(arr_x, dropout_rate, arr_y, seed)

  1197.     print(arr_y.asnumpy())

  1198.     print(seed)

  1199. 

  1200. 

  1201. def test_dropout2d():

  1202.     ctx = ht.gpu(0)

  1203.     shapeX = (2, 2, 2, 4)

  1204.     x = np.random.uniform(0, 10, size=shapeX).astype(np.float32)

  1205.     arr_x = ht.array(x, ctx=ctx)

  1206.     arr_y = ht.empty(shapeX, ctx=ctx)

  1207.     dropout_rate = 0.6

  1208.     import ctypes

  1209.     seed = ctypes.c_ulonglong(0)

  1210.     print(x)

  1211.     gpu_op.dropout2d(arr_x, dropout_rate, arr_y, seed)

  1212.     print(arr_y.asnumpy())

  1213.     print(seed)

  1214.     gpu_op.dropout2d_gradient(arr_x, dropout_rate, arr_y, seed)

  1215.     print(arr_y.asnumpy())

  1216.     print(seed)

  1217. 

  1218. 

  1219. def test_instance_norm2d():

  1220.     from hetu import ndarray

  1221.     ctx = ht.gpu(0)

  1222.     shapeX = (2, 2, 2, 4)

  1223.     eps = 0.00000001

  1224.     x = ht.array(np.arange(32).reshape(shapeX), ctx=ctx)

  1225.     mid_shape = (2, 2, 1, 1)

  1226.     mean = ndarray.empty(mid_shape, ctx=ndarray.gpu(0))

  1227.     var = ndarray.empty(mid_shape, ctx=ndarray.gpu(0))

  1228.     out = ht.empty(shapeX, ctx=ctx)

  1229.     gpu_op.instance_normalization2d(x, mean, var, out, eps)

  1230. 

  1231.     def correct_res(arr_in, eps):

  1232.         mu = np.mean(arr_in, axis=(2, 3))

  1233.         sigma = np.var(arr_in, axis=(2, 3))

  1234. 

  1235.         arr_out = np.ones_like(arr_in)

  1236.         for i in range(arr_in.shape[0]):

  1237.             for j in range(arr_in.shape[1]):

  1238.                 for k in range(arr_in.shape[2]):

  1239.                     for h in range(arr_in.shape[3]):

  1240.                         arr_out[i][j][k][h] = (

  1241.                             arr_in[i][j][k][h] - mu[i][j]) / np.sqrt(sigma[i][j] + eps)

  1242.         return arr_out

  1243. 

  1244.     np.testing.assert_allclose(

  1245.         out.asnumpy(), correct_res(x.asnumpy(), eps), rtol=1e-5)

  1246. 

  1247. 

  1248. def test_instance_norm2d_gradient():

  1249.     from hetu import ndarray

  1250.     ctx = ht.gpu(0)

  1251.     shapeX = (2, 2, 2, 4)

  1252.     eps = 0.00000001

  1253.     x = ht.array(np.arange(32).reshape(shapeX), ctx=ctx)

  1254.     mid_shape = (2, 2, 1, 1)

  1255.     mean = ndarray.empty(mid_shape, ctx=ndarray.gpu(0))

  1256.     var = ndarray.empty(mid_shape, ctx=ndarray.gpu(0))

  1257.     norm_out = ht.empty(shapeX, ctx=ctx)

  1258.     gpu_op.instance_normalization2d(x, mean, var, norm_out, eps)

  1259. 

  1260.     x_grad = ht.array(np.ones(shapeX), ctx=ctx)

  1261.     norm_grad_out = ht.empty(shapeX, ctx=ctx)

  1262.     gpu_op.instance_normalization2d_gradient(

  1263.         x_grad, x, norm_grad_out, mean, var, eps)

  1264. 

  1265.     def correct_res(arr_in, grad_in, eps):

  1266.         mu = np.mean(arr_in, axis=(2, 3))

  1267.         sigma = np.var(arr_in, axis=(2, 3))

  1268. 

  1269.         grad_out = np.ones_like(grad_in)

  1270.         for i in range(arr_in.shape[0]):

  1271.             for j in range(arr_in.shape[1]):

  1272.                 for k in range(arr_in.shape[2]):

  1273.                     for h in range(arr_in.shape[3]):

  1274.                         y = (arr_in[i][j][k][h] - mu[i][j]) / \

  1275.                             np.sqrt(sigma[i][j] + eps)

  1276.                         grad_out[i][j][k][h] = grad_in[i][j][k][h] * \

  1277.                             (1 - 1/(shapeX[2]*shapeX[3]) - y **

  1278.                              2) / np.sqrt(sigma[i][j] + eps)

  1279.         return grad_out

  1280. 

  1281.     np.testing.assert_allclose(norm_grad_out.asnumpy(), correct_res(

  1282.         x.asnumpy(), x_grad.asnumpy(), eps), rtol=1e-5)

  1283. 

  1284. 

  1285. def test_onehot():

  1286.     ctx = ht.gpu(0)

  1287.     shapein = (2, 3)

  1288.     num_classes = 7

  1289.     x = np.random.randint(7, size=shapein)

  1290.     arr_x = ht.array(x, ctx=ctx)

  1291.     arr_y = ht.empty(list(shapein) + [num_classes], ctx=ctx)

  1292.     gpu_op.one_hot(arr_x, arr_y)

  1293.     print(x)

  1294.     print(arr_y.asnumpy())

  1295. 

  1296. 

  1297. test_array_set()

  1298. test_broadcast_to()

  1299. test_reduce_sum_axis_zero()

  1300. test_matrix_elementwise_add()

  1301. test_matrix_elementwise_add_by_const()

  1302. test_matrix_elementwise_multiply()

  1303. test_matrix_elementwise_multiply_by_const()

  1304. test_matrix_elementwise_divide()

  1305. test_matrix_elementwise_divide_const()

  1306. test_matrix_opposite()

  1307. test_matrix_multiply()

  1308. test_matrix_sqrt()

  1309. test_matrix_rsqrt()

  1310. test_relu()

  1311. test_leaky_relu()

  1312. test_relu_gradient()

  1313. test_leaky_relu_gradient()

  1314. test_softmax()

  1315. test_softmax_cross_entropy()

  1316. test_softmax_cross_entropy_gradient()

  1317. test_conv2d()

  1318. test_conv2d_Gradient()

  1319. test_average_pooling()

  1320. test_reshape()

  1321. test_conv2d_broadcast_to()

  1322. test_conv2d_reduce_sum()

  1323. test_cudnn_conv2d()

  1324. test_cudnn_conv2d_Gradient()

  1325. test_CuDNN_max_pooling()

  1326. test_CuDNN_dropout_op()

  1327. test_pad()

  1328. test_concat()

  1329. test_matrix_transpose()

  1330. test_slice()

  1331. test_where()

  1332. test_batch_matrix_multiply()

  1333. test_broadcast_shape()

  1334. test_reduce_sum()

  1335. test_reduce_mean()

  1336. test_dropout()

  1337. test_dropout2d()

  1338. test_instance_norm2d()

  1339. test_instance_norm2d_gradient()

  1340. test_onehot()