upsample ok

include/Layer.h

@@ -319,6 +319,7 @@ public:
     virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
 
     int stride;
+    bool reverse;
 };
 
 struct box {

include/kernels.h

@@ -7,6 +7,8 @@ void activationELUForward(dnnType* srcData, dnnType* dstData, int size, cudaStre
 void activationLEAKYForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream = cudaStream_t(0));
 void activationLOGISTICForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream = cudaStream_t(0));
 
+void fill(dnnType* data, int size, dnnType val, cudaStream_t stream = cudaStream_t(0));
+
 void reorgForward(  dnnType* srcData, dnnType* dstData, 
                     int n, int c, int h, int w, int stride, cudaStream_t stream = cudaStream_t(0));
 void softmaxForward(float *input, int n, int batch, int batch_offset, 
@@ -17,5 +19,9 @@ void shortcutForward(dnnType* srcData, dnnType* dstData, int n1, int c1, int h1,
                                                          int n2, int c2, int h2, int w2, int s2, 
                      cudaStream_t stream = cudaStream_t(0));
 
+void upsampleForward(dnnType* srcData, dnnType* dstData, 
+                     int n, int c, int h, int w, int s, int forward, float scale,                                   
+                     cudaStream_t stream = cudaStream_t(0));
+
 void float2half(float* srcData, __half* dstData, int size, const cudaStream_t stream = cudaStream_t(0));
 #endif //KERNELS_H

src/Upsample.cpp

@@ -8,14 +8,14 @@ namespace tk { namespace dnn {
 Upsample::Upsample(Network *net, int stride) : Layer(net) {
 
     this->stride = stride;
-    
+
     output_dim.n = input_dim.n;
-    output_dim.c = input_dim.c*stride*stride;
-    output_dim.h = input_dim.h/stride;
-    output_dim.w = input_dim.w/stride;
+    output_dim.c = input_dim.c;
+    output_dim.h = input_dim.h*stride;
+    output_dim.w = input_dim.w*stride;
     output_dim.l = input_dim.l;
     
-    checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(dnnType)) );
+    checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) );
 }
 
 Upsample::~Upsample() {
@@ -25,6 +25,8 @@ Upsample::~Upsample() {
 
 dnnType* Upsample::infer(dataDim_t &dim, dnnType* srcData) {
 
+    fill(dstData, output_dim.tot(), 0.0);
+    upsampleForward(srcData, dstData, input_dim.n, input_dim.c, input_dim.h, input_dim.w, stride, 1, 1);
 
     dim = output_dim;
     return dstData;

src/kernels/fill.cu

@@ -0,0 +1,21 @@
+#include "kernels.h"
+
+__global__
+void fill_kernel(dnnType *data, int size, dnnType val) {
+
+    int i = blockDim.x*blockIdx.x + threadIdx.x;
+
+    if(i<size) {    
+        data[i] = val;
+    }
+ }
+
+
+void fill(dnnType* data, int size, dnnType val, cudaStream_t stream)
+{
+    int blocks = (size+255)/256;
+    int threads = 256;
+    fill_kernel<<<blocks, threads, 0, stream>>>(data, size, val);
+}
+
+

src/kernels/upsample.cu

@@ -0,0 +1,35 @@
+#include "kernels.h"
+
+__global__ void upsample_kernel(size_t N, dnnType *x, int w, int h, int c, int batch, int stride, int forward, float scale, dnnType *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]);
+}
+
+void upsampleForward(dnnType* srcData, dnnType* dstData, 
+                     int n, int c, int h, int w, int s, int forward, float scale,                                   
+                     cudaStream_t stream) {
+
+    int size = w*h*c*n*s*s;
+    int blocks = (size+255)/256;
+    int threads = 256;
+    upsample_kernel<<<blocks, threads, 0, stream>>>(size, srcData, w, h, c, n, s, forward, scale, dstData);
+}

tests/yolo3_berkeley/yolo3_berkeley.cpp

@@ -62,7 +62,7 @@ const char *c79_bin     = "../tests/yolo3_berkeley/layers/c79.bin";
 const char *c80_bin     = "../tests/yolo3_berkeley/layers/c80.bin";
 const char *c81_bin     = "../tests/yolo3_berkeley/layers/c81.bin";
 const char *c84_bin     = "../tests/yolo3_berkeley/layers/c84.bin";
-const char *output_bin = "../tests/yolo3_berkeley/debug/layer84_out.bin";
+const char *output_bin = "../tests/yolo3_berkeley/debug/layer85_out.bin";
 
 int main() {
 
@@ -225,6 +225,7 @@ int main() {
     tk::dnn::Route      m83  (&net, m83_layers, 1);
     tk::dnn::Conv2d     c84  (&net, 256, 1, 1, 1, 1, 0, 0, c84_bin, true);
     tk::dnn::Activation a84  (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Upsample   u85  (&net, 2);
 
 
     // Load input