From 77841359a9890abcda0edab3bb269c2123d1684b Mon Sep 17 00:00:00 2001 From: liuqi Date: Mon, 25 Dec 2017 17:56:16 +0800 Subject: [PATCH] Use mul24 to replace mul operator. --- mace/kernels/opencl/cl/space_to_batch.cl | 22 +++++++++++----------- 1 file changed, 11 insertions(+), 11 deletions(-) diff --git a/mace/kernels/opencl/cl/space_to_batch.cl b/mace/kernels/opencl/cl/space_to_batch.cl index 2725eae0..9ad63509 100644 --- a/mace/kernels/opencl/cl/space_to_batch.cl +++ b/mace/kernels/opencl/cl/space_to_batch.cl @@ -17,25 +17,25 @@ __kernel void space_to_batch(__read_only image2d_t space_data, const int batch_b_idx = batch_hb_idx / batch_height; const int batch_h_idx = batch_hb_idx % batch_height; - const int block_size = block_height * block_width; + const int block_size = mul24(block_height, block_width); const int space_b_idx = batch_b_idx / block_size; const int remaining_batch_idx = batch_b_idx % block_size; const int space_h_idx = (remaining_batch_idx / block_width) + - batch_h_idx * block_height - padding_height; + mul24(batch_h_idx, block_height) - padding_height; const int space_w_idx = (remaining_batch_idx % block_width) + - batch_w_idx * block_width - padding_width; + mul24(batch_w_idx, block_width) - padding_width; - const int space_coord_x = select(chan_idx * space_width + space_w_idx, + const int space_coord_x = select(mul24(chan_idx, space_width) + space_w_idx, -1, space_w_idx < 0 || space_w_idx >= space_width); - const int space_coord_y = select(space_b_idx * space_height + space_h_idx, + const int space_coord_y = select(mul24(space_b_idx, space_height) + space_h_idx, -1, space_h_idx < 0 || space_h_idx >= space_height); int2 space_coord = (int2)(space_coord_x, space_coord_y); DATA_TYPE4 value = READ_IMAGET(space_data, SAMPLER, space_coord); - int2 batch_coord = (int2)(chan_idx * batch_width + batch_w_idx, batch_hb_idx); + int2 batch_coord = (int2)(mul24(chan_idx, batch_width) + batch_w_idx, batch_hb_idx); WRITE_IMAGET(batch_data, batch_coord, value); } @@ -56,20 +56,20 @@ __kernel void batch_to_space(__read_only image2d_t batch_data, const int batch_b_idx = batch_hb_idx / batch_height; const int batch_h_idx = batch_hb_idx % batch_height; - const int block_size = block_height * block_width; + const int block_size = mul24(block_height, block_width); const int space_b_idx = batch_b_idx / block_size; const int remaining_batch_idx = batch_b_idx % block_size; const int space_h_idx = (remaining_batch_idx / block_width) + - batch_h_idx * block_height - padding_height; + mul24(batch_h_idx, block_height) - padding_height; const int space_w_idx = (remaining_batch_idx % block_width) + - batch_w_idx * block_width - padding_width; + mul24(batch_w_idx, block_width) - padding_width; if (0 <= space_w_idx && space_w_idx < space_width && 0 <= space_h_idx && space_h_idx < space_height) { - int2 batch_coord = (int2)(chan_idx * batch_width + batch_w_idx, batch_hb_idx); + int2 batch_coord = (int2)(mul24(chan_idx, batch_width) + batch_w_idx, batch_hb_idx); DATA_TYPE4 value = READ_IMAGET(batch_data, SAMPLER, batch_coord); - int2 space_coord = (int2)(chan_idx * space_width + space_w_idx, + int2 space_coord = (int2)(mul24(chan_idx, space_width) + space_w_idx, space_b_idx * space_height + space_h_idx); WRITE_IMAGET(space_data, space_coord, value); } -- GitLab