From a392bbd0c957a00e3782c96e7ced84a29ff9dd88 Mon Sep 17 00:00:00 2001
From: Joseph Redmon <pjreddie@gmail.com>
Date: Tue, 15 Mar 2016 05:33:02 +0000
Subject: [PATCH] Play along w/ alphago
---
src/blas_kernels.cu | 95 ++++++++++++++++++++++++++++++++++++-----------
1 files changed, 73 insertions(+), 22 deletions(-)
diff --git a/src/blas_kernels.cu b/src/blas_kernels.cu
index 8f05eb9..98366f8 100644
--- a/src/blas_kernels.cu
+++ b/src/blas_kernels.cu
@@ -1,6 +1,7 @@
#include "cuda_runtime.h"
#include "curand.h"
#include "cublas_v2.h"
+#include <assert.h>
extern "C" {
#include "blas.h"
@@ -14,7 +15,7 @@
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)
@@ -23,7 +24,7 @@
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)
@@ -45,7 +46,7 @@
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)
@@ -82,7 +83,7 @@
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));
}
}
@@ -110,7 +111,7 @@
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.));
}
}
@@ -127,7 +128,7 @@
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)
@@ -166,7 +167,7 @@
__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;
@@ -287,7 +288,7 @@
for(i = 0; i < threads; ++i){
variance[filter] += local[i];
}
- variance[filter] /= spatial * batch;
+ variance[filter] /= (spatial * batch - 1);
}
}
@@ -374,26 +375,76 @@
check_error(cudaPeekAtLastError());
}
-__global__ void shortcut_kernel(int size, float *out, int w, int h, int c, int batch, int sample, float *add, int stride, int c2, int min_c)
+__global__ void shortcut_kernel(int size, int minw, int minh, int minc, int stride, int sample, int batch, int w1, int h1, int c1, float *add, int w2, int h2, int c2, float *out)
{
int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
if (id >= size) return;
- int i = id % (w/sample);
- id /= (w/sample);
- int j = id % (h/sample);
- id /= (h/sample);
- int k = id % min_c;
- id /= min_c;
- int b = id;
- int out_index = i*sample + w*(j*sample + h*(k + c*b));
- int add_index = b*w*stride/sample*h*stride/sample*c2 + i*stride + w*stride/sample*(j*stride + h*stride/sample*k);
+ int i = id % minw;
+ id /= minw;
+ int j = id % minh;
+ id /= minh;
+ int k = id % minc;
+ id /= minc;
+ int b = id % batch;
+
+ int out_index = i*sample + w2*(j*sample + h2*(k + c2*b));
+ int add_index = i*stride + w1*(j*stride + h1*(k + c1*b));
out[out_index] += add[add_index];
}
-extern "C" void shortcut_gpu(float *out, int w, int h, int c, int batch, int sample, float *add, int stride, int c2)
+extern "C" void shortcut_gpu(int batch, int w1, int h1, int c1, float *add, int w2, int h2, int c2, float *out)
{
- int min_c = (c < c2) ? c : c2;
- int size = batch * w/sample * h/sample * min_c;
- shortcut_kernel<<<cuda_gridsize(size), BLOCK>>>(size, out, w, h, c, batch, sample, add, stride, c2, min_c);
+ int minw = (w1 < w2) ? w1 : w2;
+ int minh = (h1 < h2) ? h1 : h2;
+ int minc = (c1 < c2) ? c1 : c2;
+
+ int stride = w1/w2;
+ int sample = w2/w1;
+ assert(stride == h1/h2);
+ assert(sample == h2/h1);
+ if(stride < 1) stride = 1;
+ if(sample < 1) sample = 1;
+
+ int size = batch * minw * minh * minc;
+ shortcut_kernel<<<cuda_gridsize(size), BLOCK>>>(size, minw, minh, minc, stride, sample, batch, w1, h1, c1, add, w2, h2, c2, out);
+ check_error(cudaPeekAtLastError());
+}
+
+__global__ void smooth_l1_kernel(int n, float *pred, float *truth, float *delta, float *error)
+{
+ int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+ if(i < n){
+ float diff = truth[i] - pred[i];
+ float abs_val = abs(diff);
+ if(abs_val < 1) {
+ error[i] = diff * diff;
+ delta[i] = diff;
+ }
+ else {
+ error[i] = 2*abs_val - 1;
+ delta[i] = (diff < 0) ? -1 : 1;
+ }
+ }
+}
+
+extern "C" void smooth_l1_gpu(int n, float *pred, float *truth, float *delta, float *error)
+{
+ smooth_l1_kernel<<<cuda_gridsize(n), BLOCK>>>(n, pred, truth, delta, error);
+ check_error(cudaPeekAtLastError());
+}
+
+__global__ void l2_kernel(int n, float *pred, float *truth, float *delta, float *error)
+{
+ int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+ if(i < n){
+ float diff = truth[i] - pred[i];
+ error[i] = diff * diff; //I know this is technically wrong, deal with it.
+ delta[i] = diff;
+ }
+}
+
+extern "C" void l2_gpu(int n, float *pred, float *truth, float *delta, float *error)
+{
+ l2_kernel<<<cuda_gridsize(n), BLOCK>>>(n, pred, truth, delta, error);
check_error(cudaPeekAtLastError());
}
--
Gitblit v1.10.0