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

Merge commit '04bd5d7ff46270c732b76cd9656e22de930e138e' into tree

Browse Source
This commit is contained in:
Francesco Gatti
2020-04-28 14:54:18 +02:00
parent 3ba276a236 04bd5d7ff4
commit e6d435a463
22 changed files with 183 additions and 144 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
### BatchSize bigger than 1
```
export TKDNN_BATCHSIZE=2
```
## 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();
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;
}
};
scripts/test_all_tests.sh
+28 -65
View File
@@ -32,6 +32,14 @@ function print_output {
out_file=results.log
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 2 ) # FP32 and FP16
# modes=( 1 2 3 ) # FP32, FP16 and INT8
@@ -57,73 +65,28 @@ do
echo -e "${ORANGE}Test INT8${NC}"
fi
export TKDNN_BATCHSIZE=2
echo -e "${ORANGE}Batch $TKDNN_BATCHSIZE ${NC}"
./test_imuodom &>> $out_file
res_imuodom=$?
print_output $res_imuodom test_imuodom
print_output $? imuodom
./test_resnet101_cnet &>> $out_file
res_resnet101_cnet=$?
print_output $res_resnet101_cnet test_resnet101_cnet
./test_yolo3 &>> $out_file
res_yolo3=$?
print_output $res_yolo3 test_yolo3
./test_yolo3_flir &>> $out_file
res_yolo3_flir=$?
print_output $res_yolo3_flir test_yolo3_flir
./test_yolo3_512 &>> $out_file
res_yolo3_512=$?
print_output $res_yolo3_512 test_yolo3_512
./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
test_net resnet101_cnet
test_net yolo3
test_net yolo3_flir
test_net yolo3_512
test_net yolo3_tiny
test_net csresnext50-panet-spp
test_net mobilenetv2ssd
test_net yolo3_tiny512
test_net yolo_tiny
test_net mobilenetv2ssd512
test_net mnist
test_net yolo
test_net yolo3_berkeley
test_net yolo_voc
test_net dla34_cnet
test_net yolo3_coco4
done
src/DeformConv2d.cpp
+2 -2
View File
@@ -102,13 +102,13 @@ dnnType* DeformConv2d::infer(dataDim_t &dim, dnnType* srcData) {
dcnV2CudaForward(stat, handle,
srcData, this->data_d,
this->bias2_d, ones_d1,
offset, mask,
offset, mask,
dstData, ones_d2,
this->kernelH, this->kernelW,
this->strideH, this->strideW,
this->paddingH, this->paddingW,
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,
this->output_dim.n, this->output_dim.c, this->output_dim.h, this->output_dim.w,
chunk_dim);
src/Network.cpp
+4
View File
@@ -34,6 +34,10 @@ Network::Network(dataDim_t input_dim) {
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"))
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;
dtRT = DataType::kFLOAT;
builderRT->setMaxBatchSize(1);
builderRT->setMaxBatchSize(net->maxBatchSize);
builderRT->setMaxWorkspaceSize(1 << 30);
if(net->fp16 && builderRT->platformHasFastFp16()) {
@@ -133,6 +133,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
input->setName("out");
networkRT->markOutput(*input);
std::cout<<"Selected maxBatchSize: "<<builderRT->getMaxBatchSize()<<"\n";
std::cout<<"Building tensorRT cuda engine...\n";
#if NV_TENSORRT_MAJOR >= 6
engineRT = builderRT->buildEngineWithConfig(*networkRT, *configRT);
@@ -181,9 +182,11 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
// create GPU buffers and a stream
for(int i=0; i<engineRT->getNbBindings(); 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));
}
@@ -192,19 +195,24 @@ NetworkRT::~NetworkRT() {
}
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));
contextRT->enqueue(1, buffersRT, stream, nullptr);
checkCuda(cudaMemcpyAsync(output, buffersRT[buf_output_idx], output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
cudaStreamSynchronize(stream);
checkCuda(cudaMemcpyAsync(buffersRT[buf_input_idx], data, batches*input_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
contextRT->enqueue(batches, buffersRT, stream, nullptr);
checkCuda(cudaMemcpyAsync(output, buffersRT[buf_output_idx], batches*output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
checkCuda(cudaStreamSynchronize(stream));
dim = output_dim;
dim.n = batches;
return output;
}
void NetworkRT::enqueue() {
contextRT->enqueue(1, buffersRT, stream, nullptr);
void NetworkRT::enqueue(int batchSize) {
contextRT->enqueue(batchSize, buffersRT, stream, nullptr);
}
ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) {
@@ -320,7 +328,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) {
lRT = (ILayer*) lRTconv;
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);
@@ -527,7 +535,7 @@ 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);
@@ -535,7 +543,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *l) {
inputs[0] = input;
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,
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);
@@ -562,7 +570,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *l) {
Weights power{dtRT, power_b, l->outputs};
Weights shift{dtRT, mean_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,
shift, scale, power);
@@ -622,7 +630,7 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
const char * buf = reinterpret_cast<const char*>(serialData);
std::string name(layerName);
std::cout<<name<<std::endl;
//std::cout<<name<<std::endl;
if(name.find("ActivationLeaky") == 0) {
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 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 chunk_dim, cudaStream_t stream)
{
// 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 height = in_h;
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,
n, m, k, &alpha,
ones, k, bias, k,
&beta, output, n);
&beta, output + batch * out_c * out_h * out_w, n);
if (stat != CUBLAS_STATUS_SUCCESS)
FatalError("CUBLAS initialization failed\n");
modulatedDeformableIm2colCuda(stream,
input, offset,
mask,
input + batch * channels * height * width,
offset,// + b * 2 * int((float)chunk_dim / batch),
mask,// + b * int((float)chunk_dim / batch),
1, channels, height, width,
height_out, width_out, deformable_group, columns);
// modulatedDeformableIm2colCudaGeneralVersion(stream,
@@ -290,7 +292,7 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle,
stat = cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N,
n, m, k, &alpha,
columns, n, weight, k,
&beta, output, n);
&beta, output + batch * out_c * out_h * out_w, n);
if (stat != CUBLAS_STATUS_SUCCESS)
FatalError("CUBLAS initialization failed\n");
tests/test_rtinference/rtinference.cpp
+40 -13
View File
@@ -2,33 +2,60 @@
#include "tkdnn.h"
#include <stdlib.h> /* srand, rand */
int main(int argc, char *argv[]) {
if(argc < 2 || !fileExist(argv[1]))
FatalError("unable to read serialRT file");
int BATCH_SIZE = 1;
if(argc >2)
BATCH_SIZE = atoi(argv[2]);
//always same test
srand (0);
//convert network to tensorRT
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()];
dnnType *output = new float[netRT.input_dim.tot()];
int ret_tensorrt = 0;
std::cout<<"Testing with batchsize: "<<BATCH_SIZE<<"\n";
printCenteredTitle(" TENSORRT inference ", '=', 30);
for(int i=0; i<100; i++) {
for(int j=0; j<netRT.input_dim.tot(); j++)
input[j] = ((float) rand() / (RAND_MAX));
for(int i=0; i<10; i++) {
// 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
checkCuda( cudaMemcpyAsync(netRT.buffersRT[netRT.buf_input_idx], input,
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);
netRT.infer(dim, input_d);
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;
}
tests/yolo3_512/yolo3_512.cpp
+1 -1
View File
@@ -10,7 +10,7 @@ int main() {
// create yolo3 model
std::string bin_path = "yolo3_512";
downloadWeightsifDoNotExist("yolo3_512/layers/input.bin", bin_path, "https://cloud.hipert.unimore.it/s/e7HfScx77JEHeYb/download");
downloadWeightsifDoNotExist("yolo3_512/layers/input.bin", bin_path, "https://cloud.hipert.unimore.it/s/RGecMeGLD4cXEWL/download");
int classes = 80;
tk::dnn::Yolo *yolo [3];
#include "models/Yolo3.h"