| | |
| | | if (index >= N) return; |
| | | int f = (index/spatial)%filters; |
| | | |
| | | x[index] = (x[index] - mean[f])/(sqrt(variance[f]) + .00001f); |
| | | x[index] = (x[index] - mean[f])/(sqrt(variance[f]) + .000001f); |
| | | } |
| | | |
| | | __global__ void normalize_delta_kernel(int N, float *x, float *mean, float *variance, float *mean_delta, float *variance_delta, int batch, int filters, int spatial, float *delta) |
| | |
| | | if (index >= N) return; |
| | | int f = (index/spatial)%filters; |
| | | |
| | | delta[index] = delta[index] * 1./(sqrt(variance[f]) + .00001f) + variance_delta[f] * 2. * (x[index] - mean[f]) / (spatial * batch) + mean_delta[f]/(spatial*batch); |
| | | delta[index] = delta[index] * 1./(sqrt(variance[f]) + .000001f) + variance_delta[f] * 2. * (x[index] - mean[f]) / (spatial * batch) + mean_delta[f]/(spatial*batch); |
| | | } |
| | | |
| | | extern "C" void normalize_delta_gpu(float *x, float *mean, float *variance, float *mean_delta, float *variance_delta, int batch, int filters, int spatial, float *delta) |
| | |
| | | variance_delta[i] += delta[index]*(x[index] - mean[i]); |
| | | } |
| | | } |
| | | variance_delta[i] *= -.5 * pow(variance[i] + .00001f, (float)(-3./2.)); |
| | | variance_delta[i] *= -.5 * pow(variance[i] + .000001f, (float)(-3./2.)); |
| | | } |
| | | |
| | | __global__ void accumulate_kernel(float *x, int n, int groups, float *sum) |
| | |
| | | for(i = 0; i < threads; ++i){ |
| | | mean_delta[filter] += local[i]; |
| | | } |
| | | mean_delta[filter] *= (-1./sqrt(variance[filter] + .00001f)); |
| | | mean_delta[filter] *= (-1./sqrt(variance[filter] + .000001f)); |
| | | } |
| | | } |
| | | |
| | |
| | | for(i = 0; i < threads; ++i){ |
| | | variance_delta[filter] += local[i]; |
| | | } |
| | | variance_delta[filter] *= -.5 * pow(variance[filter] + .00001f, (float)(-3./2.)); |
| | | variance_delta[filter] *= -.5 * pow(variance[filter] + .000001f, (float)(-3./2.)); |
| | | } |
| | | } |
| | | |
| | |
| | | mean_delta[i] += delta[index]; |
| | | } |
| | | } |
| | | mean_delta[i] *= (-1./sqrt(variance[i] + .00001f)); |
| | | mean_delta[i] *= (-1./sqrt(variance[i] + .000001f)); |
| | | } |
| | | |
| | | extern "C" void mean_delta_gpu(float *delta, float *variance, int batch, int filters, int spatial, float *mean_delta) |
| | |
| | | |
| | | __global__ void variance_kernel(float *x, float *mean, int batch, int filters, int spatial, float *variance) |
| | | { |
| | | float scale = 1./(batch * spatial); |
| | | float scale = 1./(batch * spatial - 1); |
| | | int j,k; |
| | | int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x; |
| | | if (i >= filters) return; |
| | |
| | | for(i = 0; i < threads; ++i){ |
| | | variance[filter] += local[i]; |
| | | } |
| | | variance[filter] /= spatial * batch; |
| | | variance[filter] /= (spatial * batch - 1); |
| | | } |
| | | } |
| | | |