   downsample_kernel<<<grids, threads>>>(
       output.getData(),
       x.getData(),
-      scale, H, W);
+      (uint8_t) scale, (uint8_t) H, (uint8_t) W);
 
   CCE(cudaDeviceSynchronize());
 

    return;

  float sum = 0; 
  for (int i = 0; i < kW; ++i)
    for(int j = 0; j < kH; ++j)
      sum += K[ i * kH + j ] * D[ (tx + i) * HEIGHT_STEP + (ty + j) ]; 

  output[ x * fH + y ] = sum;
}

      HEIGHT_STEP = blockDim.y + kH - 1;

  int nTotal = WIDTH_STEP * HEIGHT_STEP;

  for (; tid<nTotal; tid += nThreads) {
    int xx = tid / HEIGHT_STEP + x0,
	yy = tid % HEIGHT_STEP + y0;

    if (xx >= W || yy >= H) continue;

    // rotate data 180 degree
    D[tid] = data[ (W - 1 - xx) * H + (H - 1 - yy) ];
  }
  __syncthreads();

  if (x >= vW || y >= vH)
    return;

  float sum = 0;
  D += tx * HEIGHT_STEP + ty;
  for (int i = 0; i < kW; ++i)
    for(int j = 0; j < kH; ++j)
      sum += K[ i * kH + j ] * D[ i * HEIGHT_STEP + j ]; 

  output[ x * vH + y ] += sum;
} 

__global__ void convn_valid_kernel_with_shm(float *output, const float *data,
