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

merge tree

Browse Source
This commit is contained in:
Francesco Gatti
2020-05-01 15:39:00 +02:00
parent c975a467b1 e6d435a463
commit 98537624cf
49 changed files with 1002 additions and 961 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/tkDNN/Network.h
+1
View File
@@ -62,6 +62,7 @@ public:
dataDim_t getOutputDim();
bool fp16, dla, int8;
int maxBatchSize;
bool dontLoadWeights;
std::string fileImgList;
std::string fileLabelList;
include/tkDNN/NetworkRT.h
+16 -1
View File
@@ -63,6 +63,7 @@ public:
const static int MAX_BUFFERS_RT = 10;
void* buffersRT[MAX_BUFFERS_RT];
dataDim_t buffersDIM[MAX_BUFFERS_RT];
int buf_input_idx, buf_output_idx;
dataDim_t input_dim, output_dim;
@@ -74,11 +75,25 @@ public:
NetworkRT(Network *net, const char *name);
virtual ~NetworkRT();
int getMaxBatchSize() {
if(engineRT != nullptr)
return engineRT->getMaxBatchSize();
else
return 0;
}
int getBuffersN() {
if(engineRT != nullptr)
return engineRT->getNbBindings();
else
return 0;
}
/**
Do inferece
*/
dnnType* infer(dataDim_t &dim, dnnType* data);
void enqueue();
void enqueue(int batchSize = 1);
nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Layer *l);
nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Conv2d *l);
include/tkDNN/kernels.h
+1 -1
View File
@@ -41,7 +41,7 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle,
const int stride_h, const int stride_w,
const int pad_h, const int pad_w,
const int dilation_h, const int dilation_w,
const int deformable_group,
const int deformable_group, const int batch_id,
const int in_n, const int in_c, const int in_h, const int in_w,
const int out_n, const int out_c, const int out_h, const int out_w,
const int dst_dim, cudaStream_t stream = cudaStream_t(0));
include/tkDNN/pluginsRT/ActivationLeakyRT.h
+1 -1
View File
@@ -42,7 +42,7 @@ public:
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
activationLEAKYForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
reinterpret_cast<dnnType*>(outputs[0]), size, stream);
reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, stream);
return 0;
}
include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h
+1 -1
View File
@@ -41,7 +41,7 @@ public:
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
activationReLUCeilingForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
reinterpret_cast<dnnType*>(outputs[0]), size, ceiling, stream);
reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, ceiling, stream);
return 0;
}
include/tkDNN/pluginsRT/ActivationSigmoidRT.h
+1 -1
View File
@@ -42,7 +42,7 @@ public:
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
activationSIGMOIDForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
reinterpret_cast<dnnType*>(outputs[0]), size, stream);
reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, stream);
return 0;
}
include/tkDNN/pluginsRT/DeformableConvRT.h
+25 -20
View File
@@ -86,26 +86,26 @@ public:
dnnType *output_conv = (dnnType*)reinterpret_cast<const dnnType*>(inputs[1]);
// split conv2d outputs into offset to mask
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(mask, mask, chunk_dim);
// deformable convolution
dcnV2CudaForward(stat, handle,
srcData, data_d,
bias2_d, ones_d1,
offset, mask,
reinterpret_cast<dnnType*>(outputs[0]), ones_d2,
kh, kw,
sh, sw,
ph, pw,
1, 1,
deformableGroup,
i_n, i_c, i_h, i_w,
o_n, o_c, o_h, o_w,
chunk_dim);
for(int b=0; b<batchSize; b++) {
checkCuda(cudaMemcpy(offset, output_conv + b * 3 * chunk_dim, 2*chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
checkCuda(cudaMemcpy(mask, output_conv + b * 3 * chunk_dim + 2*chunk_dim, chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
// kernel sigmoide
activationSIGMOIDForward(mask, mask, chunk_dim);
// deformable convolution
dcnV2CudaForward(stat, handle,
srcData, data_d,
bias2_d, ones_d1,
offset, mask,
reinterpret_cast<dnnType*>(outputs[0]), ones_d2,
kh, kw,
sh, sw,
ph, pw,
1, 1,
deformableGroup, b,
i_n, i_c, i_h, i_w,
o_n, o_c, o_h, o_w,
chunk_dim);
}
return 0;
}
@@ -185,6 +185,11 @@ public:
dnnType * offset;
dnnType * mask;
dnnType *ones_d2;
// dnnType *input_n;
// dnnType *offset_n;
// dnnType *mask_n;
// dnnType *output_n;
tk::dnn::DeformConv2d *defRT;
};
include/tkDNN/pluginsRT/FlattenConcatRT.h
+9 -5
View File
@@ -47,11 +47,15 @@ public:
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
checkCuda( cudaMemcpy(dstData, srcData, rows*cols*sizeof(dnnType), cudaMemcpyDeviceToDevice));
float const alpha(1.0);
float const beta(0.0);
checkERROR( cublasSgeam( handle, CUBLAS_OP_T, CUBLAS_OP_N, rows, cols, &alpha, srcData, cols, &beta, srcData, rows, dstData, rows ));
checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*rows*cols*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
checkERROR( cublasSetStream(handle, stream) );
for(int i=0; i<batchSize; i++) {
float const alpha(1.0);
float const beta(0.0);
int offset = i*rows*cols;
checkERROR( cublasSgeam( handle, CUBLAS_OP_T, CUBLAS_OP_N, rows, cols, &alpha, srcData + offset, cols, &beta, srcData + offset, rows, dstData + offset, rows ));
}
return 0;
}
include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h
+2 -2
View File
@@ -42,10 +42,10 @@ public:
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
std::cout<<this->n<<" "<<this->c<<" "<<this->h<<" "<<this->w<<" "<<this->stride_H<<" "<<this->stride_W<<" "<<this->winSize<<" "<<this->padding<<std::endl;
//std::cout<<this->n<<" "<<this->c<<" "<<this->h<<" "<<this->w<<" "<<this->stride_H<<" "<<this->stride_W<<" "<<this->winSize<<" "<<this->padding<<std::endl;
dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
MaxPoolingForward(srcData, dstData, this->n, this->c, this->h, this->w, this->stride_H, this->stride_W, this->winSize, this->padding);
MaxPoolingForward(srcData, dstData, batchSize, this->c, this->h, this->w, this->stride_H, this->stride_W, this->winSize, this->padding, stream);
return 0;
}
include/tkDNN/pluginsRT/RegionRT.h
+6 -6
View File
@@ -50,18 +50,18 @@ public:
for (int b = 0; b < batchSize; ++b){
for(int n = 0; n < num; ++n){
int index = entry_index(b, n*w*h, 0, batchSize);
int index = entry_index(b, n*w*h, 0);
activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream);
index = entry_index(b, n*w*h, coords, batchSize);
index = entry_index(b, n*w*h, coords);
activationLOGISTICForward(srcData + index, dstData + index, w*h, stream);
}
}
//softmax start
int index = entry_index(0, 0, coords + 1, batchSize);
int index = entry_index(0, 0, coords + 1);
softmaxForward( srcData + index, classes, batchSize*num,
(batchSize*c*h*w)/num,
(c*h*w)/num,
w*h, 1, w*h, 1, dstData + index, stream);
return 0;
@@ -85,10 +85,10 @@ public:
int c, h, w;
int classes, coords, num;
int entry_index(int batch, int location, int entry, int batchSize) {
int entry_index(int batch, int location, int entry) {
int n = location / (w*h);
int loc = location % (w*h);
return batch*c*h*w*batchSize + n*w*h*(coords+classes+1) + entry*w*h + loc;
return batch*c*h*w + n*w*h*(coords+classes+1) + entry*w*h + loc;
}
};
include/tkDNN/pluginsRT/ReshapeRT.h
+1 -1
View File
@@ -40,7 +40,7 @@ public:
dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
checkCuda( cudaMemcpy(dstData, srcData, c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice));
checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
return 0;
}
include/tkDNN/pluginsRT/RouteRT.h
+4
View File
@@ -3,6 +3,10 @@
class RouteRT : public IPlugin {
/**
THIS IS NOT USED ANYMORE
*/
public:
RouteRT() {
}
include/tkDNN/pluginsRT/ShortcutRT.h
+2 -1
View File
@@ -47,7 +47,8 @@ public:
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
shortcutForward(srcDataBack, dstData, batchSize, c, h, w, 1, batchSize, bc, bh, bw, 1, stream);
for(int b=0; b < batchSize; ++b)
shortcutForward(srcDataBack + b*bc*bh*bw, dstData + b*c*h*w, 1, c, h, w, 1, 1, bc, bh, bw, 1, stream);
return 0;
}
include/tkDNN/pluginsRT/YoloRT.h
+4 -4
View File
@@ -62,10 +62,10 @@ public:
for (int b = 0; b < batchSize; ++b){
for(int n = 0; n < n_masks; ++n){
int index = entry_index(b, n*w*h, 0, batchSize);
int index = entry_index(b, n*w*h, 0);
activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream);
index = entry_index(b, n*w*h, 4, batchSize);
index = entry_index(b, n*w*h, 4);
activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*w*h, stream);
}
}
@@ -109,10 +109,10 @@ public:
dnnType *mask;
dnnType *bias;
int entry_index(int batch, int location, int entry, int batchSize) {
int entry_index(int batch, int location, int entry) {
int n = location / (w*h);
int loc = location % (w*h);
return batch*c*h*w*batchSize + n*w*h*(4+classes+1) + entry*w*h + loc;
return batch*c*h*w + n*w*h*(4+classes+1) + entry*w*h + loc;
}
};