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:
Micaela Verucchi
2020-04-08 11:20:56 +02:00
parent 6a08118e3d 5be71cf9b7
commit b070bb3bd8
8 changed files with 92 additions and 152 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/tkDNN/Int8BatchStream.h
+12 -6
View File
@@ -19,16 +19,21 @@
#include "utils.h"
#include "tkdnn.h"
class BatchStream
{
/*
* BatchStream implements the stream for the INT8 calibrator.
* It reads the two files .txt with the list of image file names
* and the list of label file names.
* It then iterates on images and labels.
*/
class BatchStream {
public:
BatchStream(tk::dnn::dataDim_t dim, int batchSize, int maxBatches, const std::string& fileimglist, const std::string& filelabellist);
virtual ~BatchStream() {}
virtual ~BatchStream() { }
void reset(int firstBatch);
bool next();
void skip(int skipCount);
float *getBatch() { return mBatch.data();}
float *getLabels() { return mLabels.data();}
float *getBatch() { return mBatch.data(); }
float *getLabels() { return mLabels.data(); }
int getBatchesRead() const { return mBatchCount; }
int getBatchSize() const { return mBatchSize; }
nvinfer1::DimsNCHW getDims() const { return mDims; }
@@ -43,7 +48,8 @@ private:
int mBatchSize{ 0 };
int mMaxBatches{ 0 };
int mBatchCount{ 0 };
int mFileCount{ 0 }, mFileBatchPos{ 0 };
int mFileCount{ 0 };
int mFileBatchPos{ 0 };
int mImageSize{ 0 };
nvinfer1::DimsNCHW mDims;
include/tkDNN/Int8Calibrator.h
+11 -3
View File
@@ -18,9 +18,17 @@
#include "tkdnn.h"
#include "utils.h"
class Int8EntropyCalibrator : public nvinfer1::IInt8EntropyCalibrator{
/*
* Int8EntropyCalibrator implements the INT8 calibrator to achieve the
* INT8 quantization. It uses a BatchStream stream to scroll through
* images data. It also implements the calibration cache, a way to
* save the calibration process results to reduce the running time:
* the calibration process takes a long time.
*/
class Int8EntropyCalibrator : public nvinfer1::IInt8EntropyCalibrator {
public:
Int8EntropyCalibrator(BatchStream& stream, int firstBatch, const std::string& calibTableFilePath, const std::string& inputBlobName, bool readCache = true);
Int8EntropyCalibrator(BatchStream& stream, int firstBatch, const std::string& calibTableFilePath,
const std::string& inputBlobName, bool readCache = true);
virtual ~Int8EntropyCalibrator() { checkCuda(cudaFree(mDeviceInput)); }
int getBatchSize() const override { return mStream.getBatchSize(); }
bool getBatch(void* bindings[], const char* names[], int nbBindings) override;
@@ -29,7 +37,7 @@ public:
private:
BatchStream mStream;
const std::string mCalibTableFilePath{nullptr};
const std::string mCalibTableFilePath{ nullptr };
const std::string mInputBlobName;
bool mReadCache{ true };
include/tkDNN/kernels.h
+1 -14
View File
@@ -32,20 +32,7 @@ void upsampleForward(dnnType *srcData, dnnType *dstData,
void float2half(float *srcData, __half *dstData, int size, const cudaStream_t stream = cudaStream_t(0));
// void modulated_deformable_im2col_cuda(cudaStream_t stream,
// const float *data_im, const float *data_offset, const float *data_mask,
// const int batch_size, const int channels, const int height_im, const int width_im,
// const int height_col, const int width_col, const int kernel_h, const int kenerl_w,
// const int pad_h, const int pad_w, const int stride_h, const int stride_w,
// const int dilation_h, const int dilation_w,
// const int deformable_group, float *data_col);
void modulated_deformable_im2col_cuda(cudaStream_t stream,
const float *data_im, const float *data_offset, const float *data_mask,
const int batch_size, const int channels, const int height_im, const int width_im,
const int height_col, const int width_col,
const int deformable_group, float *data_col);
void dcn_v2_cuda_forward(cublasStatus_t stat, cublasHandle_t handle,
void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle,
float *input, float *weight,
float *bias, float *ones,
float *offset, float *mask,
include/tkDNN/pluginsRT/DeformableConvRT.h
+6 -26
View File
@@ -12,12 +12,6 @@ public:
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;
@@ -53,13 +47,11 @@ public:
checkCuda( cudaMemcpy(ones_d2, deformable->ones_d2, sizeof(dnnType)*dim_ones, cudaMemcpyDeviceToDevice) );
}
stat = cublasCreate(&handle);
if (stat != CUBLAS_STATUS_SUCCESS) {
printf ("CUBLAS initialization failed\n");
return;
}
if (stat != CUBLAS_STATUS_SUCCESS)
FatalError("CUBLAS initialization failed\n");
}
~DeformableConvRT(){
~DeformableConvRT() {
checkCuda( cudaFree(data_d) );
checkCuda( cudaFree(bias2_d) );
checkCuda( cudaFree(ones_d1) );
@@ -77,24 +69,13 @@ public:
return DimsCHW{defRT->output_dim.c, defRT->output_dim.h, defRT->output_dim.w};
}
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];
}
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { }
int initialize() override {
return 0;
}
virtual void terminate() override {
}
virtual void terminate() override { }
virtual size_t getWorkspaceSize(int maxBatchSize) const override {
return 0;
@@ -111,7 +92,7 @@ public:
activationSIGMOIDForward(mask, mask, chunk_dim);
// deformable convolution
dcn_v2_cuda_forward(stat, handle,
dcnV2CudaForward(stat, handle,
srcData, data_d,
bias2_d, ones_d1,
offset, mask,
@@ -205,6 +186,5 @@ public:
dnnType * mask;
dnnType *ones_d2;
tk::dnn::DeformConv2d *defRT;
};