Minor fix

AlexeyAB

2018-05-27 160eddddc4e265d5ee59a38797c30720bf46cd7c

 src/blas_kernels.cu

@@ -784,3 +784,34 @@
    check_error(cudaPeekAtLastError());
}


__global__ void upsample_kernel(size_t N, float *x, int w, int h, int c, int batch, int stride, int forward, float scale, float *out)
{
   size_t i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
   if (i >= N) return;
   int out_index = i;
   int out_w = i % (w*stride);
   i = i / (w*stride);
   int out_h = i % (h*stride);
   i = i / (h*stride);
   int out_c = i%c;
   i = i / c;
   int b = i%batch;

   int in_w = out_w / stride;
   int in_h = out_h / stride;
   int in_c = out_c;

   int in_index = b*w*h*c + in_c*w*h + in_h*w + in_w;


   if (forward) out[out_index] += scale * x[in_index];
   else atomicAdd(x + in_index, scale * out[out_index]);
}

extern "C" void upsample_gpu(float *in, int w, int h, int c, int batch, int stride, int forward, float scale, float *out)
{
   size_t size = w*h*c*batch*stride*stride;
   upsample_kernel << <cuda_gridsize(size), BLOCK >> >(size, in, w, h, c, batch, stride, forward, scale, out);
   check_error(cudaPeekAtLastError());
}