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 ae1d8cd9e6 0e5c90634f
commit 5f6011b206
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
+9 -3
View File
@@ -19,8 +19,13 @@
#include "utils.h" #include "utils.h"
#include "tkdnn.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: public:
BatchStream(tk::dnn::dataDim_t dim, int batchSize, int maxBatches, const std::string& fileimglist, const std::string& filelabellist); BatchStream(tk::dnn::dataDim_t dim, int batchSize, int maxBatches, const std::string& fileimglist, const std::string& filelabellist);
virtual ~BatchStream() { } virtual ~BatchStream() { }
@@ -43,7 +48,8 @@ private:
int mBatchSize{ 0 }; int mBatchSize{ 0 };
int mMaxBatches{ 0 }; int mMaxBatches{ 0 };
int mBatchCount{ 0 }; int mBatchCount{ 0 };
int mFileCount{ 0 }, mFileBatchPos{ 0 }; int mFileCount{ 0 };
int mFileBatchPos{ 0 };
int mImageSize{ 0 }; int mImageSize{ 0 };
nvinfer1::DimsNCHW mDims; nvinfer1::DimsNCHW mDims;
include/tkDNN/Int8Calibrator.h
+9 -1
View File
@@ -18,9 +18,17 @@
#include "tkdnn.h" #include "tkdnn.h"
#include "utils.h" #include "utils.h"
/*
* 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 { class Int8EntropyCalibrator : public nvinfer1::IInt8EntropyCalibrator {
public: 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)); } virtual ~Int8EntropyCalibrator() { checkCuda(cudaFree(mDeviceInput)); }
int getBatchSize() const override { return mStream.getBatchSize(); } int getBatchSize() const override { return mStream.getBatchSize(); }
bool getBatch(void* bindings[], const char* names[], int nbBindings) override; bool getBatch(void* bindings[], const char* names[], int nbBindings) override;
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 float2half(float *srcData, __half *dstData, int size, const cudaStream_t stream = cudaStream_t(0));
// void modulated_deformable_im2col_cuda(cudaStream_t stream, void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle,
// 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,
float *input, float *weight, float *input, float *weight,
float *bias, float *ones, float *bias, float *ones,
float *offset, float *mask, float *offset, float *mask,
include/tkDNN/pluginsRT/DeformableConvRT.h
+5 -25
View File
@@ -12,12 +12,6 @@ public:
int o_n, int o_c, int o_h, int o_w, int o_n, int o_c, int o_h, int o_w,
tk::dnn::DeformConv2d *deformable = nullptr) { tk::dnn::DeformConv2d *deformable = nullptr) {
this->chunk_dim = chunk_dim; 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->kh = kh;
this->kw = kw; this->kw = kw;
this->sh = sh; this->sh = sh;
@@ -53,10 +47,8 @@ public:
checkCuda( cudaMemcpy(ones_d2, deformable->ones_d2, sizeof(dnnType)*dim_ones, cudaMemcpyDeviceToDevice) ); checkCuda( cudaMemcpy(ones_d2, deformable->ones_d2, sizeof(dnnType)*dim_ones, cudaMemcpyDeviceToDevice) );
} }
stat = cublasCreate(&handle); stat = cublasCreate(&handle);
if (stat != CUBLAS_STATUS_SUCCESS) { if (stat != CUBLAS_STATUS_SUCCESS)
printf ("CUBLAS initialization failed\n"); FatalError("CUBLAS initialization failed\n");
return;
}
} }
~DeformableConvRT() { ~DeformableConvRT() {
@@ -77,24 +69,13 @@ public:
return DimsCHW{defRT->output_dim.c, defRT->output_dim.h, defRT->output_dim.w}; 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 { 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 { int initialize() override {
return 0; return 0;
} }
virtual void terminate() override { virtual void terminate() override { }
}
virtual size_t getWorkspaceSize(int maxBatchSize) const override { virtual size_t getWorkspaceSize(int maxBatchSize) const override {
return 0; return 0;
@@ -111,7 +92,7 @@ public:
activationSIGMOIDForward(mask, mask, chunk_dim); activationSIGMOIDForward(mask, mask, chunk_dim);
// deformable convolution // deformable convolution
dcn_v2_cuda_forward(stat, handle, dcnV2CudaForward(stat, handle,
srcData, data_d, srcData, data_d,
bias2_d, ones_d1, bias2_d, ones_d1,
offset, mask, offset, mask,
@@ -205,6 +186,5 @@ public:
dnnType * mask; dnnType * mask;
dnnType *ones_d2; dnnType *ones_d2;
tk::dnn::DeformConv2d *defRT; tk::dnn::DeformConv2d *defRT;
}; };
src/DeformConv2d.cpp
+16 -20
View File
@@ -10,10 +10,9 @@ namespace tk { namespace dnn {
void DeformConv2d::initCUDNN() { void DeformConv2d::initCUDNN() {
stat = cublasCreate(&handle); stat = cublasCreate(&handle);
if (stat != CUBLAS_STATUS_SUCCESS) { if (stat != CUBLAS_STATUS_SUCCESS)
printf ("CUBLAS initialization failed\n"); FatalError("CUBLAS initialization failed\n");
return;
}
checkCUDNN( cudnnCreateTensorDescriptor(&biasTensorDesc) ); checkCUDNN( cudnnCreateTensorDescriptor(&biasTensorDesc) );
checkCUDNN( cudnnSetTensor4dDescriptor(biasTensorDesc, checkCUDNN( cudnnSetTensor4dDescriptor(biasTensorDesc,
net->tensorFormat, net->dataType, net->tensorFormat, net->dataType,
@@ -28,27 +27,26 @@ void DeformConv2d::initCUDNN() {
int dst_dim = preconv->output_dim.tot(); int dst_dim = preconv->output_dim.tot();
if( dst_dim % 3 != 0 ) if( dst_dim % 3 != 0 )
std::cout<<"take attention\n\n"; FatalError("DeformConv2d: the Conv2d output is not divisible by three");
chunk_dim = dst_dim/3; chunk_dim = dst_dim/3;
checkCuda( cudaMalloc(&offset, 2*chunk_dim*sizeof(dnnType))); checkCuda( cudaMalloc(&offset, 2*chunk_dim*sizeof(dnnType)));
checkCuda( cudaMalloc(&mask, chunk_dim*sizeof(dnnType))); checkCuda( cudaMalloc(&mask, chunk_dim*sizeof(dnnType)));
// kernel ones // kernel ones
checkCuda( cudaMalloc(&ones_d1, (height_ones*width_ones)*sizeof(dnnType)) ); checkCuda( cudaMalloc(&ones_d1, (height_ones*width_ones)*sizeof(dnnType)) );
dnnType *aus1; dnnType *ones_h1;
checkCuda( cudaMallocHost(&aus1, (height_ones*width_ones)*sizeof(dnnType)) ); checkCuda( cudaMallocHost(&ones_h1, (height_ones*width_ones)*sizeof(dnnType)) );
for(int i=0; i<height_ones*width_ones; i++) for(int i=0; i<height_ones*width_ones; i++)
aus1[i]=1.0f; ones_h1[i]=1.0f;
checkCuda( cudaMemcpy(ones_d1, aus1, (height_ones*width_ones)*sizeof(dnnType), cudaMemcpyHostToDevice) ); checkCuda( cudaMemcpy(ones_d1, ones_h1, (height_ones*width_ones)*sizeof(dnnType), cudaMemcpyHostToDevice) );
checkCuda( cudaFreeHost(aus1) ); checkCuda( cudaFreeHost(ones_h1) );
checkCuda( cudaMalloc(&ones_d2, dim_ones*sizeof(dnnType)) ); checkCuda( cudaMalloc(&ones_d2, dim_ones*sizeof(dnnType)) );
dnnType *aus2; dnnType *ones_h2;
checkCuda( cudaMallocHost(&aus2, dim_ones*sizeof(dnnType)) ); checkCuda( cudaMallocHost(&ones_h2, dim_ones*sizeof(dnnType)) );
for(int i=0; i<dim_ones; i++) for(int i=0; i<dim_ones; i++)
aus2[i]=1.0f; ones_h2[i]=1.0f;
checkCuda( cudaMemcpy(ones_d2, aus2, (dim_ones)*sizeof(dnnType), cudaMemcpyHostToDevice) ); checkCuda( cudaMemcpy(ones_d2, ones_h2, (dim_ones)*sizeof(dnnType), cudaMemcpyHostToDevice) );
checkCuda( cudaFreeHost(aus2) ); checkCuda( cudaFreeHost(ones_h2) );
checkCuda( cudaDeviceSynchronize() ); checkCuda( cudaDeviceSynchronize() );
} }
@@ -58,7 +56,6 @@ DeformConv2d::DeformConv2d( Network *net, int out_ch, int deformable_group, int
LayerWgs(net, net->getOutputDim().c, out_ch, kernelH, kernelW, 1, LayerWgs(net, net->getOutputDim().c, out_ch, kernelH, kernelW, 1,
d_fname_weights, batchnorm, true) { d_fname_weights, batchnorm, true) {
this->out_ch = out_ch; this->out_ch = out_ch;
this->deformableGroup = deformable_group; this->deformableGroup = deformable_group;
this->kernelH = kernelH; this->kernelH = kernelH;
@@ -81,7 +78,6 @@ DeformConv2d::DeformConv2d( Network *net, int out_ch, int deformable_group, int
} }
DeformConv2d::~DeformConv2d() { DeformConv2d::~DeformConv2d() {
checkCUDNN( cudnnDestroyTensorDescriptor(biasTensorDesc) ); checkCUDNN( cudnnDestroyTensorDescriptor(biasTensorDesc) );
checkCuda( cudaFree(dstData) ); checkCuda( cudaFree(dstData) );
checkCuda( cudaFree(ones_d1) ); checkCuda( cudaFree(ones_d1) );
@@ -96,14 +92,14 @@ dnnType* DeformConv2d::infer(dataDim_t &dim, dnnType* srcData) {
// conv2d // conv2d
output_conv = preconv->infer(dim, srcData); output_conv = preconv->infer(dim, srcData);
// split conv2d outputs into offset to mask // split conv2d outputs into offset and mask
checkCuda(cudaMemcpy(offset, output_conv, 2*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)); checkCuda(cudaMemcpy(mask, output_conv + 2*chunk_dim, chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
// kernel sigmoide // kernel sigmoide
activationSIGMOIDForward(mask, mask, chunk_dim); activationSIGMOIDForward(mask, mask, chunk_dim);
// deformable convolution // deformable convolution
dcn_v2_cuda_forward(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,
src/Int8BatchStream.cpp
+10 -23
View File
@@ -5,8 +5,7 @@
#include <opencv2/highgui/highgui.hpp> #include <opencv2/highgui/highgui.hpp>
#include <opencv2/imgproc/imgproc.hpp> #include <opencv2/imgproc/imgproc.hpp>
BatchStream::BatchStream(tk::dnn::dataDim_t dim, int batchSize, int maxBatches, const std::string& fileimglist, const std::string& filelabellist) BatchStream::BatchStream(tk::dnn::dataDim_t dim, int batchSize, int maxBatches, const std::string& fileimglist, const std::string& filelabellist) {
{
mBatchSize = batchSize; mBatchSize = batchSize;
mMaxBatches = maxBatches; mMaxBatches = maxBatches;
mDims = nvinfer1::DimsNCHW{ dim.n, dim.c, dim.h, dim.w }; mDims = nvinfer1::DimsNCHW{ dim.n, dim.c, dim.h, dim.w };
@@ -25,22 +24,19 @@ BatchStream::BatchStream(tk::dnn::dataDim_t dim, int batchSize, int maxBatches,
reset(0); reset(0);
} }
void BatchStream::reset(int firstBatch) void BatchStream::reset(int firstBatch) {
{
mBatchCount = 0; mBatchCount = 0;
mFileCount = 0; mFileCount = 0;
mFileBatchPos = mDims.n(); mFileBatchPos = mDims.n();
skip(firstBatch); skip(firstBatch);
} }
bool BatchStream::next() bool BatchStream::next() {
{
std::cout<<"Next batch: "<<mBatchCount<<" of "<<mMaxBatches<<"\n"; std::cout<<"Next batch: "<<mBatchCount<<" of "<<mMaxBatches<<"\n";
if (mBatchCount == mMaxBatches-1) if (mBatchCount == mMaxBatches-1)
return false; return false;
for (int csize = 1, batchPos = 0; batchPos < mBatchSize; batchPos += csize, mFileBatchPos += csize) for (int csize = 1, batchPos = 0; batchPos < mBatchSize; batchPos += csize, mFileBatchPos += csize) {
{
assert(mFileBatchPos > 0 && mFileBatchPos <= mDims.n()); assert(mFileBatchPos > 0 && mFileBatchPos <= mDims.n());
if (mFileBatchPos == mDims.n() && !update()) if (mFileBatchPos == mDims.n() && !update())
return false; return false;
@@ -53,10 +49,8 @@ bool BatchStream::next()
return true; return true;
} }
void BatchStream::skip(int skipCount) void BatchStream::skip(int skipCount) {
{ if (mBatchSize >= mDims.n() && mBatchSize%mDims.n() == 0 && mFileBatchPos == mDims.n()) {
if (mBatchSize >= mDims.n() && mBatchSize%mDims.n() == 0 && mFileBatchPos == mDims.n())
{
mFileCount += skipCount * mBatchSize / mDims.n(); mFileCount += skipCount * mBatchSize / mDims.n();
return; return;
} }
@@ -67,8 +61,7 @@ void BatchStream::skip(int skipCount)
mBatchCount = x; mBatchCount = x;
} }
void BatchStream::readInListFile(const std::string& dataFilePath, std::vector<std::string>& mListIn) void BatchStream::readInListFile(const std::string& dataFilePath, std::vector<std::string>& mListIn) {
{
// dataFilePath contains the list of image paths // dataFilePath contains the list of image paths
int count = 0; int count = 0;
FILE* f = fopen(dataFilePath.c_str(), "r"); FILE* f = fopen(dataFilePath.c_str(), "r");
@@ -91,8 +84,7 @@ void BatchStream::readInListFile(const std::string& dataFilePath, std::vector<st
fclose(f); fclose(f);
} }
void BatchStream::readCVimage(std::string inputFileName, std::vector<float>& res, bool fixshape) void BatchStream::readCVimage(std::string inputFileName, std::vector<float>& res, bool fixshape) {
{
// unaltered original DsImage // unaltered original DsImage
cv::Mat m_OrigImage; cv::Mat m_OrigImage;
// letterboxed DsImage given to the network as input // letterboxed DsImage given to the network as input
@@ -138,19 +130,15 @@ void BatchStream::readCVimage(std::string inputFileName, std::vector<float>& res
res.assign(m_LetterboxImage.begin<float>(), m_LetterboxImage.end<float>()); res.assign(m_LetterboxImage.begin<float>(), m_LetterboxImage.end<float>());
} }
void BatchStream::readLabels(std::string inputFileName, std::vector<float>& ris) void BatchStream::readLabels(std::string inputFileName, std::vector<float>& ris) {
{
std::ifstream is(inputFileName.c_str()); std::ifstream is(inputFileName.c_str());
//read only the first number: the image sub-portion class //read only the first number: the image sub-portion class
while (true) { while (true) {
float val; float val;
// Read
is >> val; is >> val;
// Check
if (!is) { if (!is) {
break; break;
} }
// Use
// insert the first number and skip all others // insert the first number and skip all others
ris.push_back(val); ris.push_back(val);
while( true ) { while( true ) {
@@ -162,8 +150,7 @@ void BatchStream::readLabels(std::string inputFileName, std::vector<float>& ris)
} }
} }
bool BatchStream::update() bool BatchStream::update() {
{
std::string imgFileName = mListImg[mFileCount]; std::string imgFileName = mListImg[mFileCount];
std::string labelFileName = mListLabel[mFileCount]; std::string labelFileName = mListLabel[mFileCount];
mFileCount++; mFileCount++;
src/Int8Calibrator.cpp
+4 -9
View File
@@ -7,17 +7,14 @@ Int8EntropyCalibrator::Int8EntropyCalibrator(BatchStream& stream, int firstBatch
mStream(stream), mStream(stream),
mCalibTableFilePath(calibTableFilePath), mCalibTableFilePath(calibTableFilePath),
mInputBlobName(inputBlobName.c_str()), mInputBlobName(inputBlobName.c_str()),
mReadCache(readCache) mReadCache(readCache) {
{
nvinfer1::DimsNCHW dims = mStream.getDims(); nvinfer1::DimsNCHW dims = mStream.getDims();
mInputCount = mStream.getBatchSize() * dims.c() * dims.h() * dims.w(); mInputCount = mStream.getBatchSize() * dims.c() * dims.h() * dims.w();
checkCuda(cudaMalloc(&mDeviceInput, mInputCount * sizeof(float))); checkCuda(cudaMalloc(&mDeviceInput, mInputCount * sizeof(float)));
mStream.reset(firstBatch); mStream.reset(firstBatch);
std::cout<<"mCalibTableFilePath\n";
} }
bool Int8EntropyCalibrator::getBatch(void* bindings[], const char* names[], int nbBindings) bool Int8EntropyCalibrator::getBatch(void* bindings[], const char* names[], int nbBindings) {
{
if (!mStream.next()) if (!mStream.next())
return false; return false;
@@ -27,8 +24,7 @@ bool Int8EntropyCalibrator::getBatch(void* bindings[], const char* names[], int
return true; return true;
} }
const void* Int8EntropyCalibrator::readCalibrationCache(size_t& length) const void* Int8EntropyCalibrator::readCalibrationCache(size_t& length) {
{
mCalibrationCache.clear(); mCalibrationCache.clear();
assert(!mCalibTableFilePath.empty()); assert(!mCalibTableFilePath.empty());
std::ifstream input(mCalibTableFilePath, std::ios::binary); std::ifstream input(mCalibTableFilePath, std::ios::binary);
@@ -42,8 +38,7 @@ const void* Int8EntropyCalibrator::readCalibrationCache(size_t& length)
return length ? &mCalibrationCache[0] : nullptr; return length ? &mCalibrationCache[0] : nullptr;
} }
void Int8EntropyCalibrator::writeCalibrationCache(const void* cache, size_t length) void Int8EntropyCalibrator::writeCalibrationCache(const void* cache, size_t length) {
{
assert(!mCalibTableFilePath.empty()); assert(!mCalibTableFilePath.empty());
std::ofstream output(mCalibTableFilePath, std::ios::binary); std::ofstream output(mCalibTableFilePath, std::ios::binary);
output.write(reinterpret_cast<const char*>(cache), length); output.write(reinterpret_cast<const char*>(cache), length);
src/kernels/deformable_conv.cu
+25 -44
View File
@@ -1,6 +1,8 @@
#include <cstdio> #include <cstdio>
#include <algorithm> #include <algorithm>
#include <cstring> #include <cstring>
#include <string>
#include <iostream>
#include "kernels.h" #include "kernels.h"
#include <errno.h> #include <errno.h>
@@ -17,8 +19,7 @@ inline int GET_BLOCKS(const int N)
__device__ float dmcn_im2col_bilinear(const float *bottom_data, const int data_width, __device__ float dmcn_im2col_bilinear(const float *bottom_data, const int data_width,
const int height, const int width, float h, float w) const int height, const int width, float h, float w) {
{
int h_low = floor(h); int h_low = floor(h);
int w_low = floor(w); int w_low = floor(w);
int h_high = h_low + 1; int h_high = h_low + 1;
@@ -44,8 +45,7 @@ __global__ void modulated_deformable_im2col_gpu_kernel(const int n,
const int height, const int width, const int height, const int width,
const int batch_size, const int num_channels, const int deformable_group, const int batch_size, const int num_channels, const int deformable_group,
const int height_col, const int width_col, const int height_col, const int width_col,
float *data_col) float *data_col) {
{
CUDA_KERNEL_LOOP(index, n) CUDA_KERNEL_LOOP(index, n)
{ {
//If n is a power of 2, ( i / n ) is equivalent to ( i ≫ log2 n ) and ( i % n ) is equivalent to ( i & n - 1 ). //If n is a power of 2, ( i / n ) is equivalent to ( i ≫ log2 n ) and ( i % n ) is equivalent to ( i & n - 1 ).
@@ -77,11 +77,9 @@ __global__ void modulated_deformable_im2col_gpu_kernel(const int n,
const float *data_mask_ptr = data_mask + add_ptr; const float *data_mask_ptr = data_mask + add_ptr;
#pragma unroll #pragma unroll
for (int i = 0; i < 3; ++i) for (int i = 0; i < 3; ++i) {
{
#pragma unroll #pragma unroll
for (int j = 0; j < 3; ++j) for (int j = 0; j < 3; ++j) {
{
const int iter_member = (i * 3 + j); const int iter_member = (i * 3 + j);
// const int data_offset_h_ptr = ((2 * (i * kernel_w + j)) * height_col + h_col) * width_col + w_col; // const int data_offset_h_ptr = ((2 * (i * kernel_w + j)) * height_col + h_col) * width_col + w_col;
const int data_offset_h_ptr = first_member + s_col2 * iter_member; const int data_offset_h_ptr = first_member + s_col2 * iter_member;
@@ -99,8 +97,7 @@ __global__ void modulated_deformable_im2col_gpu_kernel(const int n,
const float w_im = offset_w + w_in + j; const float w_im = offset_w + w_in + j;
//if (h_im >= 0 && w_im >= 0 && h_im < height && w_im < width) { //if (h_im >= 0 && w_im >= 0 && h_im < height && w_im < width) {
float val = static_cast<float>(0); float val = static_cast<float>(0);
if (h_im < height && w_im < width && h_im > -1 && w_im > -1) if (h_im < height && w_im < width && h_im > -1 && w_im > -1) {
{
//const float map_h = i * dilation_h + offset_h; //const float map_h = i * dilation_h + offset_h;
//const float map_w = j * dilation_w + offset_w; //const float map_w = j * dilation_w + offset_w;
//const int cur_height = height - h_in; //const int cur_height = height - h_in;
@@ -116,7 +113,7 @@ __global__ void modulated_deformable_im2col_gpu_kernel(const int n,
} }
} }
__global__ void modulated_deformable_im2col_gpu_kernel2(const int n, __global__ void modulated_deformable_im2col_gpu_kernel_general_version(const int n,
const float *data_im, const float *data_offset, const float *data_mask, const float *data_im, const float *data_offset, const float *data_mask,
const int height, const int width, const int kernel_h, const int kernel_w, const int height, const int width, const int kernel_h, const int kernel_w,
const int pad_h, const int pad_w, const int pad_h, const int pad_w,
@@ -125,12 +122,10 @@ __global__ void modulated_deformable_im2col_gpu_kernel2(const int n,
const int channel_per_deformable_group, const int channel_per_deformable_group,
const int batch_size, const int num_channels, const int deformable_group, const int batch_size, const int num_channels, const int deformable_group,
const int height_col, const int width_col, const int height_col, const int width_col,
float *data_col) float *data_col) {
{
CUDA_KERNEL_LOOP(index, n) CUDA_KERNEL_LOOP(index, n)
{ {
//If n is a power of 2, ( i / n ) is equivalent to ( i ≫ log2 n ) and ( i % n ) is equivalent to ( i & n - 1 ). //If n is a power of 2, ( i / n ) is equivalent to ( i ≫ log2 n ) and ( i % n ) is equivalent to ( i & n - 1 ).
// printf("--- %d %d %d %d %d %d %d %d\n",kernel_h, kernel_w, pad_h, pad_w, stride_h, stride_w, dilation_h, dilation_w);
const int ind_on_w = index / width_col; const int ind_on_w = index / width_col;
const int ind_on_w_on_h = ind_on_w / height_col; const int ind_on_w_on_h = ind_on_w / height_col;
const int kk = kernel_h * kernel_w; const int kk = kernel_h * kernel_w;
@@ -160,11 +155,9 @@ __global__ void modulated_deformable_im2col_gpu_kernel2(const int n,
const float *data_mask_ptr = data_mask + add_ptr; const float *data_mask_ptr = data_mask + add_ptr;
#pragma unroll #pragma unroll
for (int i = 0; i < kernel_h; ++i) for (int i = 0; i < kernel_h; ++i) {
{
#pragma unroll #pragma unroll
for (int j = 0; j < kernel_w; ++j) for (int j = 0; j < kernel_w; ++j) {
{
const int iter_member = (i * kernel_w + j); const int iter_member = (i * kernel_w + j);
// const int data_offset_h_ptr = ((2 * (i * kernel_w + j)) * height_col + h_col) * width_col + w_col; // const int data_offset_h_ptr = ((2 * (i * kernel_w + j)) * height_col + h_col) * width_col + w_col;
const int data_offset_h_ptr = first_member + s_col2 * iter_member; const int data_offset_h_ptr = first_member + s_col2 * iter_member;
@@ -182,8 +175,7 @@ __global__ void modulated_deformable_im2col_gpu_kernel2(const int n,
const float w_im = offset_w + w_in + j * dilation_w; const float w_im = offset_w + w_in + j * dilation_w;
//if (h_im >= 0 && w_im >= 0 && h_im < height && w_im < width) { //if (h_im >= 0 && w_im >= 0 && h_im < height && w_im < width) {
float val = static_cast<float>(0); float val = static_cast<float>(0);
if (h_im < height && w_im < width && h_im > -1 && w_im > -1) if (h_im < height && w_im < width && h_im > -1 && w_im > -1) {
{
//const float map_h = i * dilation_h + offset_h; //const float map_h = i * dilation_h + offset_h;
//const float map_w = j * dilation_w + offset_w; //const float map_w = j * dilation_w + offset_w;
//const int cur_height = height - h_in; //const int cur_height = height - h_in;
@@ -199,8 +191,7 @@ __global__ void modulated_deformable_im2col_gpu_kernel2(const int n,
} }
} }
void modulatedDeformableIm2colCuda(cudaStream_t stream,
void modulated_deformable_im2col_cuda(cudaStream_t stream,
const float* data_im, const float* data_offset, const float* data_mask, 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 batch_size, const int channels, const int height_im, const int width_im,
const int height_col, const int width_col, const int height_col, const int width_col,
@@ -216,13 +207,10 @@ void modulated_deformable_im2col_cuda(cudaStream_t stream,
cudaError_t err = cudaGetLastError(); cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) if (err != cudaSuccess)
{ FatalError("error in modulatedDeformableIm2colCuda: " + std::string(cudaGetErrorString(err)) + "\n");
printf("error in modulated_deformable_im2col_cuda: %s\n", cudaGetErrorString(err));
} }
} void modulatedDeformableIm2colCudaGeneralVersion(cudaStream_t stream,
void modulated_deformable_im2col_cuda2(cudaStream_t stream,
const float* data_im, const float* data_offset, const float* data_mask, 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 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 height_col, const int width_col, const int kernel_h, const int kenerl_w,
@@ -232,7 +220,7 @@ void modulated_deformable_im2col_cuda2(cudaStream_t stream,
// num_axes should be smaller than block size // num_axes should be smaller than block size
const int channel_per_deformable_group = channels / deformable_group; const int channel_per_deformable_group = channels / deformable_group;
const int num_kernels = channels * batch_size * height_col * width_col; const int num_kernels = channels * batch_size * height_col * width_col;
modulated_deformable_im2col_gpu_kernel2 modulated_deformable_im2col_gpu_kernel_general_version
<<<GET_BLOCKS(num_kernels), CUDA_NUM_THREADS, <<<GET_BLOCKS(num_kernels), CUDA_NUM_THREADS,
0, stream>>>( 0, stream>>>(
num_kernels, data_im, data_offset, data_mask, height_im, width_im, kernel_h, kenerl_w, num_kernels, data_im, data_offset, data_mask, height_im, width_im, kernel_h, kenerl_w,
@@ -241,13 +229,10 @@ void modulated_deformable_im2col_cuda2(cudaStream_t stream,
cudaError_t err = cudaGetLastError(); cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) if (err != cudaSuccess)
{ FatalError("error in modulatedDeformableIm2colCudaGeneralVersion: " + std::string(cudaGetErrorString(err)) + "\n");
printf("error in modulated_deformable_im2col_cuda: %s\n", cudaGetErrorString(err));
} }
} void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle,
void dcn_v2_cuda_forward(cublasStatus_t stat, cublasHandle_t handle,
float *input, float *weight, float *input, float *weight,
float *bias, float *ones, float *bias, float *ones,
float *offset, float *mask, float *offset, float *mask,
@@ -266,7 +251,6 @@ void dcn_v2_cuda_forward(cublasStatus_t stat, cublasHandle_t handle,
const int height = in_h; const int height = in_h;
const int width = in_w; const int width = in_w;
const int channels_out = out_c; const int channels_out = out_c;
const int height_out = (height + 2 * pad_h - (dilation_h * (kernel_h - 1) + 1)) / stride_h + 1; const int height_out = (height + 2 * pad_h - (dilation_h * (kernel_h - 1) + 1)) / stride_h + 1;
@@ -282,17 +266,15 @@ void dcn_v2_cuda_forward(cublasStatus_t stat, cublasHandle_t handle,
n, m, k, &alpha, n, m, k, &alpha,
ones, k, bias, k, ones, k, bias, k,
&beta, output, n); &beta, output, n);
if (stat != CUBLAS_STATUS_SUCCESS) { if (stat != CUBLAS_STATUS_SUCCESS)
printf ("CUBLAS initialization failed\n"); FatalError("CUBLAS initialization failed\n");
return ;
}
modulated_deformable_im2col_cuda(stream, modulatedDeformableIm2colCuda(stream,
input, offset, input, offset,
mask, mask,
1, channels, height, width, 1, channels, height, width,
height_out, width_out, deformable_group, columns); height_out, width_out, deformable_group, columns);
// modulated_deformable_im2col_cuda2(stream, // modulatedDeformableIm2colCudaGeneralVersion(stream,
// input, offset, // input, offset,
// mask, // mask,
// 1, channels, height, width, // 1, channels, height, width,
@@ -310,8 +292,7 @@ void dcn_v2_cuda_forward(cublasStatus_t stat, cublasHandle_t handle,
columns, n, weight, k, columns, n, weight, k,
&beta, output, n); &beta, output, n);
if (stat != CUBLAS_STATUS_SUCCESS) { if (stat != CUBLAS_STATUS_SUCCESS)
printf ("CUBLAS initialization failed\n"); FatalError("CUBLAS initialization failed\n");
return ;
}
} }