mmr/tkDNN
Archived
1
0
Fork 0
Code Issues 46 Pull Requests 12 Releases 4 Activity

Merge branch 'cnet' of https://github.com/ceccocats/tkDNN into cnet

Browse Source
This commit is contained in:
fbagni
2019-12-23 16:14:24 +01:00
parent 7e7f480e2b 5272f1cde6
commit 23a1365dc4
4 changed files with 188 additions and 24 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/tkDNN/Layer.h
+2
View File
@@ -43,6 +43,8 @@ public:
dataDim_t input_dim, output_dim;
dnnType *dstData; //where results will be putted
int id = 0;
bool final; //if the layer is the final one
std::string getLayerName() {
include/tkDNN/pluginsRT/DeformableConvRT.h
+135 -18
View File
@@ -7,8 +7,51 @@ class DeformableConvRT : public IPlugin {
public:
DeformableConvRT(tk::dnn::DeformConv2d *deformable) {
this->defRT = deformable;
DeformableConvRT(int chunk_dim, int kh, int kw, int sh, int sw, int ph, int pw,
int deformableGroup, int i_n, int i_c, int i_h, int i_w,
int o_n, int o_c, int o_h, int o_w,
tk::dnn::DeformConv2d *deformable = nullptr) {
this->chunk_dim = chunk_dim;
// int dst_dim = conv_dim.tot();
// std::cout<<"conv_dim: \n";
// conv_dim.print();
// if (dst_dim % 3 != 0 )
// std::cout<<"take attention\n\n";
// this->chunk_dim = dst_dim/3;
this->kh = kh;
this->kw = kw;
this->sh = sh;
this->sw = sw;
this->ph = ph;
this->pw = pw;
this->deformableGroup = deformableGroup;
this->i_n = i_n;
this->i_c = i_c;
this->i_h = i_h;
this->i_w = i_w;
this->o_n = o_n;
this->o_c = o_c;
this->o_h = o_h;
this->o_w = o_w;
height_ones = (i_h + 2 * ph - (1 * (kh - 1) + 1)) / sh + 1;
width_ones = (i_w + 2 * pw - (1 * (kw - 1) + 1)) / sw + 1;
dim_ones = i_c * kh * kw * 1 * height_ones * width_ones;
std::cout<<i_c * o_c * kh * kw * 1<<"\n";
checkCuda( cudaMalloc(&data_d, i_c * o_c * kh * kw * 1 * sizeof(dnnType)));
checkCuda( cudaMalloc(&bias2_d, o_c*sizeof(dnnType)));
checkCuda( cudaMalloc(&ones_d1, height_ones * width_ones * sizeof(dnnType)));
checkCuda( cudaMalloc(&offset, 2*chunk_dim*sizeof(dnnType)));
checkCuda( cudaMalloc(&mask, chunk_dim*sizeof(dnnType)));
checkCuda( cudaMalloc(&ones_d2, dim_ones*sizeof(dnnType)));
if(deformable != nullptr) {
this->defRT = deformable;
checkCuda( cudaMemcpy(data_d, deformable->data_d, sizeof(dnnType)*i_c * o_c * kh * kw * 1, cudaMemcpyDeviceToDevice) );
checkCuda( cudaMemcpy(bias2_d, deformable->bias2_d, sizeof(dnnType)*o_c, cudaMemcpyDeviceToDevice) );
checkCuda( cudaMemcpy(ones_d1, deformable->ones_d1, sizeof(dnnType)*height_ones*width_ones, cudaMemcpyDeviceToDevice) );
checkCuda( cudaMemcpy(offset, deformable->offset, sizeof(dnnType)*2*chunk_dim, cudaMemcpyDeviceToDevice) );
checkCuda( cudaMemcpy(mask, deformable->mask, sizeof(dnnType)*chunk_dim, cudaMemcpyDeviceToDevice) );
checkCuda( cudaMemcpy(ones_d2, deformable->ones_d2, sizeof(dnnType)*dim_ones, cudaMemcpyDeviceToDevice) );
}
}
~DeformableConvRT(){
@@ -24,6 +67,14 @@ public:
}
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
// i_n = 1;
// i_c = inputDims[0].d[0];
// i_h = inputDims[0].d[1];
// i_w = inputDims[0].d[2];
// o_n = 1;
// o_c = outputDims[0].d[0];
// o_h = outputDims[0].d[1];
// o_w = outputDims[0].d[2];
}
int initialize() override {
@@ -39,42 +90,108 @@ public:
}
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
std::cout<<"LOL\n";
dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *output_conv = (dnnType*)reinterpret_cast<const dnnType*>(inputs[1]);
// split conv2d outputs into offset to mask
checkCuda(cudaMemcpy(defRT->offset, defRT->output_conv, 2*defRT->chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
checkCuda(cudaMemcpy(defRT->mask, defRT->output_conv + 2*defRT->chunk_dim, defRT->chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
checkCuda(cudaMemcpy(offset, output_conv, 2*chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
checkCuda(cudaMemcpy(mask, output_conv + 2*chunk_dim, chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
// kernel sigmoide
activationSIGMOIDForward(defRT->mask, defRT->mask, defRT->chunk_dim);
activationSIGMOIDForward(mask, mask, chunk_dim);
// deformable convolution
dcn_v2_cuda_forward(srcData, defRT->data_d,
defRT->bias2_d, defRT->ones_d1,
defRT->offset, defRT->mask,
reinterpret_cast<dnnType*>(outputs[0]), defRT->ones_d2,
defRT->kernelH, defRT->kernelW,
defRT->strideH, defRT->strideW,
defRT->paddingH, defRT->paddingW,
dcn_v2_cuda_forward(srcData, data_d,
bias2_d, ones_d1,
offset, mask,
reinterpret_cast<dnnType*>(outputs[0]), ones_d2,
kh, kw,
sh, sw,
ph, pw,
1, 1,
defRT->deformableGroup,
defRT->preconv->input_dim.n, defRT->preconv->input_dim.c, defRT->preconv->input_dim.h, defRT->preconv->input_dim.w,
defRT->output_dim.n, defRT->output_dim.c, defRT->output_dim.h, defRT->output_dim.w,
defRT->chunk_dim);
deformableGroup,
i_n, i_c, i_h, i_w,
o_n, o_c, o_h, o_w,
chunk_dim);
return 0;
}
virtual size_t getSerializationSize() override {
return 0;
return 16 * sizeof(int) + chunk_dim * 3 * sizeof(dnnType) + (i_c * o_c * kh * kw * 1 ) * sizeof(dnnType) +
o_c * sizeof(dnnType) + height_ones * width_ones * sizeof(dnnType) + dim_ones * sizeof(dnnType);
}
virtual void serialize(void* buffer) override {
char *buf = reinterpret_cast<char*>(buffer);
tk::dnn::writeBUF(buf, chunk_dim);
tk::dnn::writeBUF(buf, kh);
tk::dnn::writeBUF(buf, kw);
tk::dnn::writeBUF(buf, sh);
tk::dnn::writeBUF(buf, sw);
tk::dnn::writeBUF(buf, ph);
tk::dnn::writeBUF(buf, pw);
tk::dnn::writeBUF(buf, deformableGroup);
tk::dnn::writeBUF(buf, i_n);
tk::dnn::writeBUF(buf, i_c);
tk::dnn::writeBUF(buf, i_h);
tk::dnn::writeBUF(buf, i_w);
tk::dnn::writeBUF(buf, o_n);
tk::dnn::writeBUF(buf, o_c);
tk::dnn::writeBUF(buf, o_h);
tk::dnn::writeBUF(buf, o_w);
dnnType *aus = new dnnType[chunk_dim*2];
checkCuda( cudaMemcpy(aus, offset, sizeof(dnnType)*2*chunk_dim, cudaMemcpyDeviceToHost) );
for(int i=0; i<chunk_dim*2; i++)
tk::dnn::writeBUF(buf, aus[i]);
free(aus);
aus = new dnnType[chunk_dim];
checkCuda( cudaMemcpy(aus, mask, sizeof(dnnType)*chunk_dim, cudaMemcpyDeviceToHost) );
for(int i=0; i<chunk_dim; i++)
tk::dnn::writeBUF(buf, aus[i]);
free(aus);
aus = new dnnType[(i_c * o_c * kh * kw * 1 )];
checkCuda( cudaMemcpy(aus, data_d, sizeof(dnnType)*(i_c * o_c * kh * kw * 1 ), cudaMemcpyDeviceToHost) );
for(int i=0; i<(i_c * o_c * kh * kw * 1 ); i++)
tk::dnn::writeBUF(buf, aus[i]);
free(aus);
aus = new dnnType[o_c];
checkCuda( cudaMemcpy(aus, bias2_d, sizeof(dnnType)*o_c, cudaMemcpyDeviceToHost) );
for(int i=0; i < o_c; i++)
tk::dnn::writeBUF(buf, aus[i]);
free(aus);
aus = new dnnType[height_ones * width_ones];
checkCuda( cudaMemcpy(aus, ones_d1, sizeof(dnnType)*height_ones * width_ones, cudaMemcpyDeviceToHost) );
for(int i=0; i<height_ones * width_ones; i++)
tk::dnn::writeBUF(buf, aus[i]);
free(aus);
aus = new dnnType[dim_ones];
checkCuda( cudaMemcpy(aus, ones_d2, sizeof(dnnType)*dim_ones, cudaMemcpyDeviceToHost) );
for(int i=0; i<dim_ones; i++)
tk::dnn::writeBUF(buf, aus[i]);
free(aus);
}
int i_n, i_c, i_h, i_w;
int o_n, o_c, o_h, o_w;
int size;
int chunk_dim;
int kh, kw;
int sh, sw;
int ph, pw;
int deformableGroup;
int height_ones;
int width_ones;
int dim_ones;
dnnType *data_d;
dnnType *bias2_d;
dnnType *ones_d1;
dnnType * offset;
dnnType * mask;
dnnType *ones_d2;
tk::dnn::DeformConv2d *defRT;
};
src/Network.cpp
+1
View File
@@ -61,6 +61,7 @@ bool Network::addLayer(Layer *l) {
if(num_layers == MAX_LAYERS)
return false;
l->id = num_layers;
layers[num_layers++] = l;
return true;
}
src/NetworkRT.cpp
+50 -6
View File
@@ -418,17 +418,21 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Upsample *l) {
}
ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *l) {
//std::cout<<"convert DEFORMABLE\n";
std::cout<<"convert DEFORMABLE\n";
ILayer *preconv = convert_layer(input, l->preconv);
checkNULL(preconv);
ITensor **inputs = new ITensor*[2];
inputs[0] = input;
inputs[1] = preconv->getOutput(0);
//std::cout<<"New plugin DEFORMABLE\n";
IPlugin *plugin = new DeformableConvRT(l);
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
std::cout<<"New plugin DEFORMABLE\n";
IPlugin *plugin = new DeformableConvRT(l->chunk_dim, l->kernelH, l->kernelW, l->strideH, l->strideW, l->paddingH, l->paddingW,
l->deformableGroup, l->input_dim.n, l->input_dim.c, l->input_dim.h, l->input_dim.w,
l->output_dim.n, l->output_dim.c, l->output_dim.h, l->output_dim.w, l);
IPluginLayer *lRT = networkRT->addPlugin(inputs, 2, *plugin);
checkNULL(lRT);
lRT->setName( ("Deformable" + std::to_string(l->id)).c_str() );
// batchnorm
void *bias_b, *power_b, *mean_b, *variance_b, *scales_b;
@@ -507,8 +511,9 @@ bool NetworkRT::deserialize(const char *filename) {
IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialData, size_t serialLength) {
const char * buf = reinterpret_cast<const char*>(serialData);
std::string name(layerName);
std::cout<<name<<std::endl;
if(name.find("Activation") == 0) {
ActivationLeakyRT *a = new ActivationLeakyRT();
@@ -578,7 +583,6 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
yolos[n_yolos++] = r;
return r;
}
if(name.find("Upsample") == 0) {
UpsampleRT *r = new UpsampleRT(readBUF<int>(buf)); //stride
r->c = readBUF<int>(buf);
@@ -598,6 +602,46 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
return r;
}
*/
if(name.find("Deformable") == 0) {
DeformableConvRT *r = new DeformableConvRT(readBUF<int>(buf), readBUF<int>(buf), readBUF<int>(buf),
readBUF<int>(buf), readBUF<int>(buf), readBUF<int>(buf),
readBUF<int>(buf), readBUF<int>(buf),
readBUF<int>(buf),readBUF<int>(buf),readBUF<int>(buf),readBUF<int>(buf),
readBUF<int>(buf),readBUF<int>(buf),readBUF<int>(buf),readBUF<int>(buf),
nullptr);
dnnType *aus = new dnnType[r->chunk_dim*2];
for(int i=0; i<r->chunk_dim*2; i++)
aus[i] = readBUF<dnnType>(buf);
checkCuda( cudaMemcpy(r->offset, aus, sizeof(dnnType)*2*r->chunk_dim, cudaMemcpyHostToDevice) );
free(aus);
aus = new dnnType[r->chunk_dim];
for(int i=0; i<r->chunk_dim; i++)
aus[i] = readBUF<dnnType>(buf);
checkCuda( cudaMemcpy(r->mask, aus, sizeof(dnnType)*r->chunk_dim, cudaMemcpyHostToDevice) );
free(aus);
aus = new dnnType[(r->i_c * r->o_c * r->kh * r->kw * 1 )];
for(int i=0; i<(r->i_c * r->o_c * r->kh * r->kw * 1 ); i++)
aus[i] = readBUF<dnnType>(buf);
checkCuda( cudaMemcpy(r->data_d, aus, sizeof(dnnType)*(r->i_c * r->o_c * r->kh * r->kw * 1 ), cudaMemcpyHostToDevice) );
free(aus);
aus = new dnnType[r->o_c];
for(int i=0; i < r->o_c; i++)
aus[i] = readBUF<dnnType>(buf);
checkCuda( cudaMemcpy(r->bias2_d, aus, sizeof(dnnType)*r->o_c, cudaMemcpyHostToDevice) );
free(aus);
aus = new dnnType[r->height_ones * r->width_ones];
for(int i=0; i<r->height_ones * r->width_ones; i++)
aus[i] = readBUF<dnnType>(buf);
checkCuda( cudaMemcpy(r->ones_d1, aus, sizeof(dnnType)*r->height_ones * r->width_ones, cudaMemcpyHostToDevice) );
free(aus);
aus = new dnnType[r->dim_ones];
for(int i=0; i<r->dim_ones; i++)
aus[i] = readBUF<dnnType>(buf);
checkCuda( cudaMemcpy(r->ones_d2, aus, sizeof(dnnType)*r->dim_ones, cudaMemcpyHostToDevice) );
free(aus);
return r;
}
FatalError("Cant deserialize Plugin");
return NULL;
}