Francesco Gatti
51 lines
1.3 KiB
Plaintext
51 lines
1.3 KiB
Plaintext
#include "kernels.h"
|
|
|
|
__global__ void reorg_kernel(int N, float *x, int w, int h, int c, int batch, int stride, int forward, float *out)
|
|
{
|
|
int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
|
|
if(i >= N) return;
|
|
int in_index = i;
|
|
int in_w = i%w;
|
|
i = i/w;
|
|
int in_h = i%h;
|
|
i = i/h;
|
|
int in_c = i%c;
|
|
i = i/c;
|
|
int b = i%batch;
|
|
|
|
int out_c = c/(stride*stride);
|
|
|
|
int c2 = in_c % out_c;
|
|
int offset = in_c / out_c;
|
|
int w2 = in_w*stride + offset % stride;
|
|
int h2 = in_h*stride + offset / stride;
|
|
//printf("%d\n", offset);
|
|
int out_index = w2 + w*stride*(h2 + h*stride*(c2 + out_c*b));
|
|
|
|
// printf("%d %d %d\n", w2, h2, c2);
|
|
//printf("%d %d\n", in_index, out_index);
|
|
//if(out_index >= N || out_index < 0) printf("bad bad bad \n");
|
|
|
|
if(forward) out[out_index] = x[in_index];
|
|
else out[in_index] = x[out_index];
|
|
//if(forward) out[1] = x[1];
|
|
//else out[0] = x[0];
|
|
}
|
|
|
|
/**
|
|
reorg function function
|
|
*/
|
|
void reorgForward(dnnType* srcData, dnnType* dstData,
|
|
int n, int c, int h, int w, int stride) {
|
|
|
|
int size = n*c*h*w;
|
|
|
|
int blocks = (size+255)/256;
|
|
int threads = 256;
|
|
|
|
reorg_kernel<<<blocks, threads>>>(size, srcData, w, h, c, n, stride, false, dstData);
|
|
checkCuda( cudaDeviceSynchronize() );
|
|
}
|
|
|
|
|