From 913d355ec1cf34aad71fdd75202fc3b0309e63a0 Mon Sep 17 00:00:00 2001
From: Joseph Redmon <pjreddie@gmail.com>
Date: Thu, 28 Jan 2016 20:30:38 +0000
Subject: [PATCH] lots of stuff
---
src/blas_kernels.cu | 16 ++++++++++++++++
1 files changed, 16 insertions(+), 0 deletions(-)
diff --git a/src/blas_kernels.cu b/src/blas_kernels.cu
index 49406db..61db29f 100644
--- a/src/blas_kernels.cu
+++ b/src/blas_kernels.cu
@@ -409,3 +409,19 @@
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