Joseph Redmon
2015-03-27 d7d7da2653ff4f79a275529b0ac3fec438880083
src/convolutional_kernels.cu
@@ -26,7 +26,7 @@
    check_error(cudaPeekAtLastError());
}
__global__ void backward_bias_kernel(float *bias_updates, float *delta, int batch, int n, int size, float scale)
__global__ void backward_bias_kernel(float *bias_updates, float *delta, int batch, int n, int size)
{
    __shared__ float part[BLOCK];
    int i,b;
@@ -42,13 +42,13 @@
    part[p] = sum;
    __syncthreads();
    if(p == 0){
        for(i = 0; i < BLOCK; ++i) bias_updates[filter] += scale * part[i];
        for(i = 0; i < BLOCK; ++i) bias_updates[filter] += part[i];
    }
}
void backward_bias_gpu(float *bias_updates, float *delta, int batch, int n, int size)
{
    backward_bias_kernel<<<n, BLOCK>>>(bias_updates, delta, batch, n, size, 1);
    backward_bias_kernel<<<n, BLOCK>>>(bias_updates, delta, batch, n, size);
    check_error(cudaPeekAtLastError());
}