Use mul24 to replace mul operator.

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);
   }