From b711627e84d2245e31a3b71d9e1119db49d6287c Mon Sep 17 00:00:00 2001
From: Joseph Redmon <pjreddie@gmail.com>
Date: Thu, 28 Jan 2016 20:30:54 +0000
Subject: [PATCH] rnn cfg
---
src/blas_kernels.cu | 56 ++++++++++++++++++++++++++++++++++++++++++--------------
1 files changed, 42 insertions(+), 14 deletions(-)
diff --git a/src/blas_kernels.cu b/src/blas_kernels.cu
index 8f05eb9..61db29f 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"
@@ -374,26 +375,53 @@
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)
+{
+ int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+ if(i < n){
+ float diff = truth[i] - pred[i];
+ if(abs(diff) > 1) delta[i] = diff;
+ else delta[i] = (diff > 0) ? 1 : -1;
+ }
+}
+
+extern "C" void smooth_l1_gpu(int n, float *pred, float *truth, float *delta)
+{
+ smooth_l1_kernel<<<cuda_gridsize(n), BLOCK>>>(n, pred, truth, delta);
check_error(cudaPeekAtLastError());
}
--
Gitblit v1.10.0