From 23955b9fa0a29465ad2a2d13c445b49e6d5adef2 Mon Sep 17 00:00:00 2001
From: Joseph Redmon <pjreddie@gmail.com>
Date: Mon, 08 Feb 2016 19:50:45 +0000
Subject: [PATCH] binary reading weights
---
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