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

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

Browse Source
This commit is contained in:
Micaela Verucchi
2020-04-28 14:43:24 +02:00
parent a761fef159 f91fb7f287
commit fdfda38fbf
21 changed files with 182 additions and 143 deletions
Download Patch File Download Diff File Expand all files Collapse all files
README.md
+5
View File
@@ -163,6 +163,11 @@ N.b. The test will be slower: this is due to the INT8 calibration, which may tak
N.b. INT8 calibration requires TensorRT version greater than or equal to 6.0 N.b. INT8 calibration requires TensorRT version greater than or equal to 6.0
### BatchSize bigger than 1
```
export TKDNN_BATCHSIZE=2
```
## mAP demo ## mAP demo
To compute mAP, precision, recall and f1score, run the map_demo. To compute mAP, precision, recall and f1score, run the map_demo.
include/tkDNN/Network.h
+1
View File
@@ -62,6 +62,7 @@ public:
dataDim_t getOutputDim(); dataDim_t getOutputDim();
bool fp16, dla, int8; bool fp16, dla, int8;
int maxBatchSize;
bool dontLoadWeights; bool dontLoadWeights;
std::string fileImgList; std::string fileImgList;
std::string fileLabelList; std::string fileLabelList;
include/tkDNN/NetworkRT.h
+16 -1
View File
@@ -63,6 +63,7 @@ public:
const static int MAX_BUFFERS_RT = 10; const static int MAX_BUFFERS_RT = 10;
void* buffersRT[MAX_BUFFERS_RT]; void* buffersRT[MAX_BUFFERS_RT];
dataDim_t buffersDIM[MAX_BUFFERS_RT];
int buf_input_idx, buf_output_idx; int buf_input_idx, buf_output_idx;
dataDim_t input_dim, output_dim; dataDim_t input_dim, output_dim;
@@ -74,11 +75,25 @@ public:
NetworkRT(Network *net, const char *name); NetworkRT(Network *net, const char *name);
virtual ~NetworkRT(); 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 Do inferece
*/ */
dnnType* infer(dataDim_t &dim, dnnType* data); 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, Layer *l);
nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Conv2d *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 stride_h, const int stride_w,
const int pad_h, const int pad_w, const int pad_h, const int pad_w,
const int dilation_h, const int dilation_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 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 out_n, const int out_c, const int out_h, const int out_w,
const int dst_dim, cudaStream_t stream = cudaStream_t(0)); 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 { 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]), 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; 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 { 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]), 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; 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 { 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]), 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; 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]); dnnType *output_conv = (dnnType*)reinterpret_cast<const dnnType*>(inputs[1]);
// split conv2d outputs into offset to mask // split conv2d outputs into offset to mask
checkCuda(cudaMemcpy(offset, output_conv, 2*chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); for(int b=0; b<batchSize; b++) {
checkCuda(cudaMemcpy(mask, output_conv + 2*chunk_dim, chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); checkCuda(cudaMemcpy(offset, output_conv + b * 3 * chunk_dim, 2*chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
// kernel sigmoide checkCuda(cudaMemcpy(mask, output_conv + b * 3 * chunk_dim + 2*chunk_dim, chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
activationSIGMOIDForward(mask, mask, chunk_dim); // kernel sigmoide
activationSIGMOIDForward(mask, mask, chunk_dim);
// deformable convolution // deformable convolution
dcnV2CudaForward(stat, handle, dcnV2CudaForward(stat, handle,
srcData, data_d, srcData, data_d,
bias2_d, ones_d1, bias2_d, ones_d1,
offset, mask, offset, mask,
reinterpret_cast<dnnType*>(outputs[0]), ones_d2, reinterpret_cast<dnnType*>(outputs[0]), ones_d2,
kh, kw, kh, kw,
sh, sw, sh, sw,
ph, pw, ph, pw,
1, 1, 1, 1,
deformableGroup, deformableGroup, b,
i_n, i_c, i_h, i_w, i_n, i_c, i_h, i_w,
o_n, o_c, o_h, o_w, o_n, o_c, o_h, o_w,
chunk_dim); chunk_dim);
}
return 0; return 0;
} }
@@ -185,6 +185,11 @@ public:
dnnType * offset; dnnType * offset;
dnnType * mask; dnnType * mask;
dnnType *ones_d2; dnnType *ones_d2;
// dnnType *input_n;
// dnnType *offset_n;
// dnnType *mask_n;
// dnnType *output_n;
tk::dnn::DeformConv2d *defRT; 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 { 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 *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]); dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
checkCuda( cudaMemcpy(dstData, srcData, rows*cols*sizeof(dnnType), cudaMemcpyDeviceToDevice)); checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*rows*cols*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
float const alpha(1.0); checkERROR( cublasSetStream(handle, stream) );
float const beta(0.0); for(int i=0; i<batchSize; i++) {
checkERROR( cublasSgeam( handle, CUBLAS_OP_T, CUBLAS_OP_N, rows, cols, &alpha, srcData, cols, &beta, srcData, rows, dstData, rows )); 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; 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 { 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 *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[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; 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 b = 0; b < batchSize; ++b){
for(int n = 0; n < num; ++n){ 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); 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); activationLOGISTICForward(srcData + index, dstData + index, w*h, stream);
} }
} }
//softmax start //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, softmaxForward( srcData + index, classes, batchSize*num,
(batchSize*c*h*w)/num, (c*h*w)/num,
w*h, 1, w*h, 1, dstData + index, stream); w*h, 1, w*h, 1, dstData + index, stream);
return 0; return 0;
@@ -85,10 +85,10 @@ public:
int c, h, w; int c, h, w;
int classes, coords, num; 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 n = location / (w*h);
int loc = 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 *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[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; return 0;
} }
include/tkDNN/pluginsRT/RouteRT.h
+4
View File
@@ -3,6 +3,10 @@
class RouteRT : public IPlugin { class RouteRT : public IPlugin {
/**
THIS IS NOT USED ANYMORE
*/
public: public:
RouteRT() { RouteRT() {
} }
include/tkDNN/pluginsRT/ShortcutRT.h
+2 -1
View File
@@ -47,7 +47,8 @@ public:
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]); dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); 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; 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 b = 0; b < batchSize; ++b){
for(int n = 0; n < n_masks; ++n){ 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); 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); activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*w*h, stream);
} }
} }
@@ -109,10 +109,10 @@ public:
dnnType *mask; dnnType *mask;
dnnType *bias; 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 n = location / (w*h);
int loc = 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;
} }
}; };
scripts/test_all_tests.sh
+28 -65
View File
@@ -32,6 +32,14 @@ function print_output {
out_file=results.log out_file=results.log
rm $out_file rm $out_file
function test_net {
./test_$1 &>> $out_file
print_output $? $1
./test_rtinference $1*.rt $TKDNN_BATCHSIZE &>> $out_file
print_output $? "batched $1"
}
modes=( 1 ) # only FP32 modes=( 1 ) # only FP32
# modes=( 1 2 ) # FP32 and FP16 # modes=( 1 2 ) # FP32 and FP16
# modes=( 1 2 3 ) # FP32, FP16 and INT8 # modes=( 1 2 3 ) # FP32, FP16 and INT8
@@ -57,73 +65,28 @@ do
echo -e "${ORANGE}Test INT8${NC}" echo -e "${ORANGE}Test INT8${NC}"
fi fi
export TKDNN_BATCHSIZE=2
echo -e "${ORANGE}Batch $TKDNN_BATCHSIZE ${NC}"
./test_imuodom &>> $out_file ./test_imuodom &>> $out_file
res_imuodom=$? print_output $? imuodom
print_output $res_imuodom test_imuodom
./test_resnet101_cnet &>> $out_file test_net resnet101_cnet
res_resnet101_cnet=$? test_net yolo3
print_output $res_resnet101_cnet test_resnet101_cnet test_net yolo3_flir
test_net yolo3_512
./test_yolo3 &>> $out_file test_net yolo3_tiny
res_yolo3=$? test_net csresnext50-panet-spp
print_output $res_yolo3 test_yolo3 test_net mobilenetv2ssd
test_net yolo3_tiny512
./test_yolo3_flir &>> $out_file test_net yolo_tiny
res_yolo3_flir=$? test_net mobilenetv2ssd512
print_output $res_yolo3_flir test_yolo3_flir test_net mnist
test_net yolo
./test_yolo3_512 &>> $out_file test_net yolo3_berkeley
res_yolo3_512=$? test_net yolo_voc
print_output $res_yolo3_512 test_yolo3_512 test_net dla34_cnet
test_net yolo3_coco4
./test_yolo3_tiny &>> $out_file
res_yolo3_tiny=$?
print_output $res_yolo3_tiny test_yolo3_tiny
./test_csresnext50-panet-spp &>> $out_file
res_csresnext50panetspp=$?
print_output $res_csresnext50panetspp "test_csresnext50-panet-spp"
./test_mobilenetv2ssd &>> $out_file
res_mobilenetv2ssd=$?
print_output $res_mobilenetv2ssd test_mobilenetv2ssd
./test_yolo3_tiny512 &>> $out_file
res_yolo3_tiny512=$?
print_output $res_yolo3_tiny512 test_yolo3_tiny512
./test_yolo_tiny &>> $out_file
res_yolo_tiny=$?
print_output $res_yolo_tiny test_yolo_tiny
./test_mobilenetv2ssd512 &>> $out_file
res_mobilenetv2ssd512=$?
print_output $res_mobilenetv2ssd512 test_mobilenetv2ssd512
./test_mnist &>> $out_file
res_mnist=$?
print_output $res_mnist test_mnist
./test_yolo &>> $out_file
res_yolo=$?
print_output $res_yolo test_yolo
./test_yolo3_berkeley &>> $out_file
res_yolo3_berkeley=$?
print_output $res_yolo3_berkeley test_yolo3_berkeley
./test_yolo_voc &>> $out_file
res_yolo_voc=$?
print_output $res_yolo_voc test_yolo_voc
./test_dla34_cnet &>> $out_file
res_dla34_cnet=$?
print_output $res_dla34_cnet test_dla34_cnet
./test_yolo3_coco4 &>> $out_file
res_yolo3_coco4=$?
print_output $res_yolo3_coco4 test_yolo3_coco4
done done
src/DeformConv2d.cpp
+2 -2
View File
@@ -102,13 +102,13 @@ dnnType* DeformConv2d::infer(dataDim_t &dim, dnnType* srcData) {
dcnV2CudaForward(stat, handle, dcnV2CudaForward(stat, handle,
srcData, this->data_d, srcData, this->data_d,
this->bias2_d, ones_d1, this->bias2_d, ones_d1,
offset, mask, offset, mask,
dstData, ones_d2, dstData, ones_d2,
this->kernelH, this->kernelW, this->kernelH, this->kernelW,
this->strideH, this->strideW, this->strideH, this->strideW,
this->paddingH, this->paddingW, this->paddingH, this->paddingW,
1, 1, 1, 1,
this->deformableGroup, this->deformableGroup, 0, //batch_id for cudnn is set to 0 (no batch)
preconv->input_dim.n, preconv->input_dim.c, preconv->input_dim.h, preconv->input_dim.w, preconv->input_dim.n, preconv->input_dim.c, preconv->input_dim.h, preconv->input_dim.w,
this->output_dim.n, this->output_dim.c, this->output_dim.h, this->output_dim.w, this->output_dim.n, this->output_dim.c, this->output_dim.h, this->output_dim.w,
chunk_dim); chunk_dim);
src/Network.cpp
+4
View File
@@ -34,6 +34,10 @@ Network::Network(dataDim_t input_dim) {
int8 = true; int8 = true;
} }
} }
maxBatchSize = 1;
if(const char* env_p = std::getenv("TKDNN_BATCHSIZE")) {
maxBatchSize = atoi(env_p);
}
if(const char* env_p = std::getenv("TKDNN_CALIB_IMG_PATH")) if(const char* env_p = std::getenv("TKDNN_CALIB_IMG_PATH"))
fileImgList = env_p; fileImgList = env_p;
src/NetworkRT.cpp
+22 -14
View File
@@ -58,7 +58,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
dataDim_t dim = net->layers[0]->input_dim; dataDim_t dim = net->layers[0]->input_dim;
dtRT = DataType::kFLOAT; dtRT = DataType::kFLOAT;
builderRT->setMaxBatchSize(1); builderRT->setMaxBatchSize(net->maxBatchSize);
builderRT->setMaxWorkspaceSize(1 << 30); builderRT->setMaxWorkspaceSize(1 << 30);
if(net->fp16 && builderRT->platformHasFastFp16()) { if(net->fp16 && builderRT->platformHasFastFp16()) {
@@ -133,6 +133,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
input->setName("out"); input->setName("out");
networkRT->markOutput(*input); networkRT->markOutput(*input);
std::cout<<"Selected maxBatchSize: "<<builderRT->getMaxBatchSize()<<"\n";
std::cout<<"Building tensorRT cuda engine...\n"; std::cout<<"Building tensorRT cuda engine...\n";
#if NV_TENSORRT_MAJOR >= 6 #if NV_TENSORRT_MAJOR >= 6
engineRT = builderRT->buildEngineWithConfig(*networkRT, *configRT); engineRT = builderRT->buildEngineWithConfig(*networkRT, *configRT);
@@ -181,9 +182,11 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
// create GPU buffers and a stream // create GPU buffers and a stream
for(int i=0; i<engineRT->getNbBindings(); i++) { for(int i=0; i<engineRT->getNbBindings(); i++) {
Dims dim = engineRT->getBindingDimensions(i); Dims dim = engineRT->getBindingDimensions(i);
checkCuda(cudaMalloc(&buffersRT[i], dim.d[0]*dim.d[1]*dim.d[2]*sizeof(dnnType))); buffersDIM[i] = dataDim_t(1, dim.d[0], dim.d[1], dim.d[2]);
std::cout<<"RtBuffer "<<i<<" dim: "; buffersDIM[i].print();
checkCuda(cudaMalloc(&buffersRT[i], engineRT->getMaxBatchSize()*dim.d[0]*dim.d[1]*dim.d[2]*sizeof(dnnType)));
} }
checkCuda(cudaMalloc(&output, output_dim.tot()*sizeof(dnnType))); checkCuda(cudaMalloc(&output, engineRT->getMaxBatchSize()*output_dim.tot()*sizeof(dnnType)));
checkCuda(cudaStreamCreate(&stream)); checkCuda(cudaStreamCreate(&stream));
} }
@@ -192,19 +195,24 @@ NetworkRT::~NetworkRT() {
} }
dnnType* NetworkRT::infer(dataDim_t &dim, dnnType* data) { dnnType* NetworkRT::infer(dataDim_t &dim, dnnType* data) {
int batches = dim.n;
if(batches > getMaxBatchSize()) {
FatalError("input batch size too large");
}
checkCuda(cudaMemcpyAsync(buffersRT[buf_input_idx], data, input_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); checkCuda(cudaMemcpyAsync(buffersRT[buf_input_idx], data, batches*input_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
contextRT->enqueue(1, buffersRT, stream, nullptr); contextRT->enqueue(batches, buffersRT, stream, nullptr);
checkCuda(cudaMemcpyAsync(output, buffersRT[buf_output_idx], output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); checkCuda(cudaMemcpyAsync(output, buffersRT[buf_output_idx], batches*output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
cudaStreamSynchronize(stream); checkCuda(cudaStreamSynchronize(stream));
dim = output_dim; dim = output_dim;
dim.n = batches;
return output; return output;
} }
void NetworkRT::enqueue() { void NetworkRT::enqueue(int batchSize) {
contextRT->enqueue(1, buffersRT, stream, nullptr); contextRT->enqueue(batchSize, buffersRT, stream, nullptr);
} }
ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) { ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) {
@@ -320,7 +328,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) {
lRT = (ILayer*) lRTconv; lRT = (ILayer*) lRTconv;
Dims d = lRTconv->getOutput(0)->getDimensions(); Dims d = lRTconv->getOutput(0)->getDimensions();
std::cout<<"DECONV: "<<d.d[0]<<" "<<d.d[1]<<" "<<d.d[2]<<" "<<d.d[3]<<"\n"; //std::cout<<"DECONV: "<<d.d[0]<<" "<<d.d[1]<<" "<<d.d[2]<<" "<<d.d[3]<<"\n";
} }
checkNULL(lRT); checkNULL(lRT);
@@ -527,7 +535,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Upsample *l) {
} }
ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *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); ILayer *preconv = convert_layer(input, l->preconv);
checkNULL(preconv); checkNULL(preconv);
@@ -535,7 +543,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *l) {
inputs[0] = input; inputs[0] = input;
inputs[1] = preconv->getOutput(0); inputs[1] = preconv->getOutput(0);
std::cout<<"New plugin DEFORMABLE\n"; //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, 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->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); l->output_dim.n, l->output_dim.c, l->output_dim.h, l->output_dim.w, l);
@@ -562,7 +570,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *l) {
Weights power{dtRT, power_b, l->outputs}; Weights power{dtRT, power_b, l->outputs};
Weights shift{dtRT, mean_b, l->outputs}; Weights shift{dtRT, mean_b, l->outputs};
Weights scale{dtRT, variance_b, l->outputs}; Weights scale{dtRT, variance_b, l->outputs};
std::cout<<lRT->getNbOutputs()<<std::endl; //std::cout<<lRT->getNbOutputs()<<std::endl;
IScaleLayer *lRT2 = networkRT->addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL, IScaleLayer *lRT2 = networkRT->addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL,
shift, scale, power); shift, scale, power);
@@ -622,7 +630,7 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
const char * buf = reinterpret_cast<const char*>(serialData); const char * buf = reinterpret_cast<const char*>(serialData);
std::string name(layerName); std::string name(layerName);
std::cout<<name<<std::endl; //std::cout<<name<<std::endl;
if(name.find("ActivationLeaky") == 0) { if(name.find("ActivationLeaky") == 0) {
ActivationLeakyRT *a = new ActivationLeakyRT(); ActivationLeakyRT *a = new ActivationLeakyRT();
src/kernels/deformable_conv.cu
+7 -5
View File
@@ -241,12 +241,13 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle,
const int stride_h, const int stride_w, const int stride_h, const int stride_w,
const int pad_h, const int pad_w, const int pad_h, const int pad_w,
const int dilation_h, const int dilation_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 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 out_n, const int out_c, const int out_h, const int out_w,
const int chunk_dim, cudaStream_t stream) const int chunk_dim, cudaStream_t stream)
{ {
// stat and handle have be moved out to preserve 2 - 6 milliseconds every 100. // stat and handle have be moved out to preserve 2 - 6 milliseconds every 100.
const int batch = batch_id;
const int channels = in_c; const int channels = in_c;
const int height = in_h; const int height = in_h;
const int width = in_w; const int width = in_w;
@@ -265,13 +266,14 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle,
stat = cublasSgemm(handle, CUBLAS_OP_T, CUBLAS_OP_N, stat = cublasSgemm(handle, CUBLAS_OP_T, CUBLAS_OP_N,
n, m, k, &alpha, n, m, k, &alpha,
ones, k, bias, k, ones, k, bias, k,
&beta, output, n); &beta, output + batch * out_c * out_h * out_w, n);
if (stat != CUBLAS_STATUS_SUCCESS) if (stat != CUBLAS_STATUS_SUCCESS)
FatalError("CUBLAS initialization failed\n"); FatalError("CUBLAS initialization failed\n");
modulatedDeformableIm2colCuda(stream, modulatedDeformableIm2colCuda(stream,
input, offset, input + batch * channels * height * width,
mask, offset,// + b * 2 * int((float)chunk_dim / batch),
mask,// + b * int((float)chunk_dim / batch),
1, channels, height, width, 1, channels, height, width,
height_out, width_out, deformable_group, columns); height_out, width_out, deformable_group, columns);
// modulatedDeformableIm2colCudaGeneralVersion(stream, // modulatedDeformableIm2colCudaGeneralVersion(stream,
@@ -290,7 +292,7 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle,
stat = cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, stat = cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N,
n, m, k, &alpha, n, m, k, &alpha,
columns, n, weight, k, columns, n, weight, k,
&beta, output, n); &beta, output + batch * out_c * out_h * out_w, n);
if (stat != CUBLAS_STATUS_SUCCESS) if (stat != CUBLAS_STATUS_SUCCESS)
FatalError("CUBLAS initialization failed\n"); FatalError("CUBLAS initialization failed\n");
tests/test_rtinference/rtinference.cpp
+40 -13
View File
@@ -2,33 +2,60 @@
#include "tkdnn.h" #include "tkdnn.h"
#include <stdlib.h> /* srand, rand */ #include <stdlib.h> /* srand, rand */
int main(int argc, char *argv[]) { int main(int argc, char *argv[]) {
if(argc < 2 || !fileExist(argv[1])) if(argc < 2 || !fileExist(argv[1]))
FatalError("unable to read serialRT file"); FatalError("unable to read serialRT file");
int BATCH_SIZE = 1;
if(argc >2)
BATCH_SIZE = atoi(argv[2]);
//always same test //always same test
srand (0); srand (0);
//convert network to tensorRT //convert network to tensorRT
tk::dnn::NetworkRT netRT(NULL, argv[1]); tk::dnn::NetworkRT netRT(NULL, argv[1]);
tk::dnn::dataDim_t idim = netRT.input_dim;
tk::dnn::dataDim_t odim = netRT.output_dim;
idim.n = BATCH_SIZE;
odim.n = BATCH_SIZE;
dnnType *input = new float[idim.tot()];
dnnType *output = new float[odim.tot()];
dnnType *input_d;
checkCuda( cudaMalloc(&input_d, idim.tot()*sizeof(dnnType)));
dnnType *input = new float[netRT.input_dim.tot()]; int ret_tensorrt = 0;
dnnType *output = new float[netRT.input_dim.tot()]; std::cout<<"Testing with batchsize: "<<BATCH_SIZE<<"\n";
printCenteredTitle(" TENSORRT inference ", '=', 30); printCenteredTitle(" TENSORRT inference ", '=', 30);
for(int i=0; i<100; i++) { for(int i=0; i<10; i++) {
for(int j=0; j<netRT.input_dim.tot(); j++)
input[j] = ((float) rand() / (RAND_MAX)); // generate input
for(int j=0; j<netRT.input_dim.tot(); j++) {
dnnType val = ((float) rand() / (RAND_MAX));
for(int b=0; b<BATCH_SIZE; b++)
input[netRT.input_dim.tot()*b + j] = val;
}
checkCuda(cudaMemcpy(input_d, input, idim.tot()*sizeof(dnnType), cudaMemcpyHostToDevice));
tk::dnn::dataDim_t dim = idim;
TIMER_START TIMER_START
checkCuda( cudaMemcpyAsync(netRT.buffersRT[netRT.buf_input_idx], input, netRT.infer(dim, input_d);
netRT.input_dim.tot()*sizeof(float), cudaMemcpyHostToDevice, netRT.stream));
netRT.enqueue();
checkCuda( cudaMemcpyAsync(output, netRT.buffersRT[netRT.buf_output_idx],
netRT.output_dim.tot()*sizeof(float), cudaMemcpyDeviceToHost, netRT.stream));
cudaStreamSynchronize(netRT.stream);
TIMER_STOP TIMER_STOP
// control output
std::cout<<"Output Buffers: "<<netRT.getBuffersN()-1<<"\n";
for(int o=1; o<netRT.getBuffersN(); o++) {
for(int b=1; b<BATCH_SIZE; b++) {
dnnType *out_d = (dnnType*) netRT.buffersRT[o];
dnnType *out0_d = out_d;
dnnType *outI_d = out_d + netRT.buffersDIM[o].tot()*b;
ret_tensorrt |= checkResult(netRT.buffersDIM[o].tot(), outI_d, out0_d) == 0 ? 0 : ERROR_TENSORRT;
}
}
} }
return 0; return ret_tensorrt;
} }