From d9ae3dd681ed1c98e807ff937dbbb9cfc4d19fe0 Mon Sep 17 00:00:00 2001
From: AlexeyAB <alexeyab84@gmail.com>
Date: Tue, 27 Mar 2018 23:59:03 +0000
Subject: [PATCH] Added Yolo v3
---
src/blas_kernels.cu | 31 +++++++++++++++++++++++++++++++
1 files changed, 31 insertions(+), 0 deletions(-)
diff --git a/src/blas_kernels.cu b/src/blas_kernels.cu
index 97b5977..1edbbbd 100644
--- a/src/blas_kernels.cu
+++ b/src/blas_kernels.cu
@@ -784,3 +784,34 @@
check_error(cudaPeekAtLastError());
}
+
+__global__ void upsample_kernel(size_t N, float *x, int w, int h, int c, int batch, int stride, int forward, float scale, float *out)
+{
+ size_t i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+ if (i >= N) return;
+ int out_index = i;
+ int out_w = i % (w*stride);
+ i = i / (w*stride);
+ int out_h = i % (h*stride);
+ i = i / (h*stride);
+ int out_c = i%c;
+ i = i / c;
+ int b = i%batch;
+
+ int in_w = out_w / stride;
+ int in_h = out_h / stride;
+ int in_c = out_c;
+
+ int in_index = b*w*h*c + in_c*w*h + in_h*w + in_w;
+
+
+ if (forward) out[out_index] += scale * x[in_index];
+ else atomicAdd(x + in_index, scale * out[out_index]);
+}
+
+extern "C" void upsample_gpu(float *in, int w, int h, int c, int batch, int stride, int forward, float scale, float *out)
+{
+ size_t size = w*h*c*batch*stride*stride;
+ upsample_kernel << <cuda_gridsize(size), BLOCK >> >(size, in, w, h, c, batch, stride, forward, scale, out);
+ check_error(cudaPeekAtLastError());
+}
\ No newline at end of file
--
Gitblit v1.10.0