From 5562f599a68b09b2cacf5434b9933b364c9f89c1 Mon Sep 17 00:00:00 2001 From: Davide Sapienza Date: Wed, 8 Apr 2020 10:22:03 +0200 Subject: [PATCH 1/2] Fix the INT8 calibrator sintax Signed-off-by: Davide Sapienza --- include/tkDNN/Int8BatchStream.h | 18 ++++++++++----- include/tkDNN/Int8Calibrator.h | 14 ++++++++--- src/Int8BatchStream.cpp | 41 +++++++++++---------------------- src/Int8Calibrator.cpp | 15 ++++-------- 4 files changed, 42 insertions(+), 46 deletions(-) diff --git a/include/tkDNN/Int8BatchStream.h b/include/tkDNN/Int8BatchStream.h index fc3863c..4349c1f 100644 --- a/include/tkDNN/Int8BatchStream.h +++ b/include/tkDNN/Int8BatchStream.h @@ -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; diff --git a/include/tkDNN/Int8Calibrator.h b/include/tkDNN/Int8Calibrator.h index 4fd4a0e..4a0ea47 100644 --- a/include/tkDNN/Int8Calibrator.h +++ b/include/tkDNN/Int8Calibrator.h @@ -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 }; diff --git a/src/Int8BatchStream.cpp b/src/Int8BatchStream.cpp index 12d04da..dd4399f 100644 --- a/src/Int8BatchStream.cpp +++ b/src/Int8BatchStream.cpp @@ -5,8 +5,7 @@ #include #include -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; mMaxBatches = maxBatches; 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); } -void BatchStream::reset(int firstBatch) -{ +void BatchStream::reset(int firstBatch) { mBatchCount = 0; mFileCount = 0; mFileBatchPos = mDims.n(); skip(firstBatch); } -bool BatchStream::next() -{ +bool BatchStream::next() { std::cout<<"Next batch: "< 0 && mFileBatchPos <= mDims.n()); if (mFileBatchPos == mDims.n() && !update()) return false; @@ -53,10 +49,8 @@ bool BatchStream::next() return true; } -void BatchStream::skip(int skipCount) -{ - if (mBatchSize >= mDims.n() && mBatchSize%mDims.n() == 0 && mFileBatchPos == mDims.n()) - { +void BatchStream::skip(int skipCount) { + if (mBatchSize >= mDims.n() && mBatchSize%mDims.n() == 0 && mFileBatchPos == mDims.n()) { mFileCount += skipCount * mBatchSize / mDims.n(); return; } @@ -67,8 +61,7 @@ void BatchStream::skip(int skipCount) mBatchCount = x; } -void BatchStream::readInListFile(const std::string& dataFilePath, std::vector& mListIn) -{ +void BatchStream::readInListFile(const std::string& dataFilePath, std::vector& mListIn) { // dataFilePath contains the list of image paths int count = 0; FILE* f = fopen(dataFilePath.c_str(), "r"); @@ -76,8 +69,8 @@ void BatchStream::readInListFile(const std::string& dataFilePath, std::vector& res, bool fixshape) -{ +void BatchStream::readCVimage(std::string inputFileName, std::vector& res, bool fixshape) { // unaltered original DsImage cv::Mat m_OrigImage; // letterboxed DsImage given to the network as input @@ -104,7 +96,7 @@ void BatchStream::readCVimage(std::string inputFileName, std::vector& res int m_Height = m_OrigImage.rows; int m_Width = m_OrigImage.cols; - if(fixshape){ + if(fixshape) { m_Height = mHeight; m_Width = mWidth; } @@ -138,22 +130,18 @@ void BatchStream::readCVimage(std::string inputFileName, std::vector& res res.assign(m_LetterboxImage.begin(), m_LetterboxImage.end()); } -void BatchStream::readLabels(std::string inputFileName, std::vector& ris) -{ +void BatchStream::readLabels(std::string inputFileName, std::vector& ris) { std::ifstream is(inputFileName.c_str()); //read only the first number: the image sub-portion class while (true) { float val; - // Read is >> val; - // Check if (!is) { break; } - // Use // insert the first number and skip all others ris.push_back(val); - while( true ){ + while( true ) { char c; is >> c; if (is.peek() == '\n') //detect "\n" @@ -162,8 +150,7 @@ void BatchStream::readLabels(std::string inputFileName, std::vector& ris) } } -bool BatchStream::update() -{ +bool BatchStream::update() { std::string imgFileName = mListImg[mFileCount]; std::string labelFileName = mListLabel[mFileCount]; mFileCount++; diff --git a/src/Int8Calibrator.cpp b/src/Int8Calibrator.cpp index 9dfd662..773a9d8 100644 --- a/src/Int8Calibrator.cpp +++ b/src/Int8Calibrator.cpp @@ -7,17 +7,14 @@ Int8EntropyCalibrator::Int8EntropyCalibrator(BatchStream& stream, int firstBatch mStream(stream), mCalibTableFilePath(calibTableFilePath), mInputBlobName(inputBlobName.c_str()), - mReadCache(readCache) -{ + mReadCache(readCache) { nvinfer1::DimsNCHW dims = mStream.getDims(); mInputCount = mStream.getBatchSize() * dims.c() * dims.h() * dims.w(); checkCuda(cudaMalloc(&mDeviceInput, mInputCount * sizeof(float))); 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()) return false; @@ -27,8 +24,7 @@ bool Int8EntropyCalibrator::getBatch(void* bindings[], const char* names[], int return true; } -const void* Int8EntropyCalibrator::readCalibrationCache(size_t& length) -{ +const void* Int8EntropyCalibrator::readCalibrationCache(size_t& length) { mCalibrationCache.clear(); assert(!mCalibTableFilePath.empty()); std::ifstream input(mCalibTableFilePath, std::ios::binary); @@ -42,10 +38,9 @@ const void* Int8EntropyCalibrator::readCalibrationCache(size_t& length) 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()); std::ofstream output(mCalibTableFilePath, std::ios::binary); output.write(reinterpret_cast(cache), length); output.close(); -} +} \ No newline at end of file From 0e5c90634fcc826f1611ff255034f198ff589195 Mon Sep 17 00:00:00 2001 From: Davide Sapienza Date: Wed, 8 Apr 2020 11:14:34 +0200 Subject: [PATCH 2/2] Fix the Deformable convolution code sintax. Signed-off-by: Davide Sapienza --- include/tkDNN/kernels.h | 15 +---- include/tkDNN/pluginsRT/DeformableConvRT.h | 32 ++-------- src/DeformConv2d.cpp | 40 ++++++------- src/kernels/deformable_conv.cu | 69 ++++++++-------------- 4 files changed, 50 insertions(+), 106 deletions(-) diff --git a/include/tkDNN/kernels.h b/include/tkDNN/kernels.h index 13fda1f..afcc168 100644 --- a/include/tkDNN/kernels.h +++ b/include/tkDNN/kernels.h @@ -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, diff --git a/include/tkDNN/pluginsRT/DeformableConvRT.h b/include/tkDNN/pluginsRT/DeformableConvRT.h index f4d0b30..4804c03 100644 --- a/include/tkDNN/pluginsRT/DeformableConvRT.h +++ b/include/tkDNN/pluginsRT/DeformableConvRT.h @@ -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; }; diff --git a/src/DeformConv2d.cpp b/src/DeformConv2d.cpp index 8cf1a6d..06bba7f 100644 --- a/src/DeformConv2d.cpp +++ b/src/DeformConv2d.cpp @@ -10,10 +10,9 @@ namespace tk { namespace dnn { void DeformConv2d::initCUDNN() { stat = cublasCreate(&handle); - if (stat != CUBLAS_STATUS_SUCCESS) { - printf ("CUBLAS initialization failed\n"); - return; - } + if (stat != CUBLAS_STATUS_SUCCESS) + FatalError("CUBLAS initialization failed\n"); + checkCUDNN( cudnnCreateTensorDescriptor(&biasTensorDesc) ); checkCUDNN( cudnnSetTensor4dDescriptor(biasTensorDesc, net->tensorFormat, net->dataType, @@ -27,28 +26,27 @@ void DeformConv2d::initCUDNN() { const int dim_ones = preconv->input_dim.c * this->kernelH * this->kernelW * 1 * height_ones * width_ones; int dst_dim = preconv->output_dim.tot(); - if (dst_dim % 3 != 0 ) - std::cout<<"take attention\n\n"; + if( dst_dim % 3 != 0 ) + FatalError("DeformConv2d: the Conv2d output is not divisible by three"); chunk_dim = dst_dim/3; checkCuda( cudaMalloc(&offset, 2*chunk_dim*sizeof(dnnType))); checkCuda( cudaMalloc(&mask, chunk_dim*sizeof(dnnType))); // kernel ones - checkCuda( cudaMalloc(&ones_d1, (height_ones*width_ones)*sizeof(dnnType)) ); - dnnType *aus1; - checkCuda( cudaMallocHost(&aus1, (height_ones*width_ones)*sizeof(dnnType)) ); + dnnType *ones_h1; + checkCuda( cudaMallocHost(&ones_h1, (height_ones*width_ones)*sizeof(dnnType)) ); for(int i=0; igetOutputDim().c, out_ch, kernelH, kernelW, 1, - d_fname_weights, batchnorm, true){ - + d_fname_weights, batchnorm, true) { this->out_ch = out_ch; this->deformableGroup = deformable_group; this->kernelH = kernelH; @@ -81,7 +78,6 @@ DeformConv2d::DeformConv2d( Network *net, int out_ch, int deformable_group, int } DeformConv2d::~DeformConv2d() { - checkCUDNN( cudnnDestroyTensorDescriptor(biasTensorDesc) ); checkCuda( cudaFree(dstData) ); checkCuda( cudaFree(ones_d1) ); @@ -96,14 +92,14 @@ dnnType* DeformConv2d::infer(dataDim_t &dim, dnnType* srcData) { // conv2d 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(mask, output_conv + 2*chunk_dim, chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); // kernel sigmoide activationSIGMOIDForward(mask, mask, chunk_dim); // deformable convolution - dcn_v2_cuda_forward(stat, handle, + dcnV2CudaForward(stat, handle, srcData, this->data_d, this->bias2_d, ones_d1, offset, mask, diff --git a/src/kernels/deformable_conv.cu b/src/kernels/deformable_conv.cu index 672655f..e46c579 100644 --- a/src/kernels/deformable_conv.cu +++ b/src/kernels/deformable_conv.cu @@ -1,6 +1,8 @@ #include #include #include +#include +#include #include "kernels.h" #include @@ -17,8 +19,7 @@ inline int GET_BLOCKS(const int N) __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 w_low = floor(w); 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 batch_size, const int num_channels, const int deformable_group, const int height_col, const int width_col, - float *data_col) -{ + float *data_col) { 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 ). @@ -77,11 +77,9 @@ __global__ void modulated_deformable_im2col_gpu_kernel(const int n, const float *data_mask_ptr = data_mask + add_ptr; #pragma unroll - for (int i = 0; i < 3; ++i) - { + for (int i = 0; i < 3; ++i) { #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 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; @@ -99,8 +97,7 @@ __global__ void modulated_deformable_im2col_gpu_kernel(const int n, const float w_im = offset_w + w_in + j; //if (h_im >= 0 && w_im >= 0 && h_im < height && w_im < width) { float val = static_cast(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_w = j * dilation_w + offset_w; //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 int height, const int width, const int kernel_h, const int kernel_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 batch_size, const int num_channels, const int deformable_group, const int height_col, const int width_col, - float *data_col) -{ + float *data_col) { 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 ). - // 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_on_h = ind_on_w / height_col; 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; #pragma unroll - for (int i = 0; i < kernel_h; ++i) - { + for (int i = 0; i < kernel_h; ++i) { #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 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; @@ -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; //if (h_im >= 0 && w_im >= 0 && h_im < height && w_im < width) { float val = static_cast(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_w = j * dilation_w + offset_w; //const int cur_height = height - h_in; @@ -199,8 +191,7 @@ __global__ void modulated_deformable_im2col_gpu_kernel2(const int n, } } - -void modulated_deformable_im2col_cuda(cudaStream_t stream, +void modulatedDeformableIm2colCuda(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, @@ -216,13 +207,10 @@ void modulated_deformable_im2col_cuda(cudaStream_t stream, cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) - { - printf("error in modulated_deformable_im2col_cuda: %s\n", cudaGetErrorString(err)); - } - + FatalError("error in modulatedDeformableIm2colCuda: " + std::string(cudaGetErrorString(err)) + "\n"); } -void modulated_deformable_im2col_cuda2(cudaStream_t stream, +void modulatedDeformableIm2colCudaGeneralVersion(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, @@ -232,7 +220,7 @@ void modulated_deformable_im2col_cuda2(cudaStream_t stream, // num_axes should be smaller than block size const int channel_per_deformable_group = channels / deformable_group; const int num_kernels = channels * batch_size * height_col * width_col; - modulated_deformable_im2col_gpu_kernel2 + modulated_deformable_im2col_gpu_kernel_general_version <<>>( 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(); if (err != cudaSuccess) - { - printf("error in modulated_deformable_im2col_cuda: %s\n", cudaGetErrorString(err)); - } - + FatalError("error in modulatedDeformableIm2colCudaGeneralVersion: " + std::string(cudaGetErrorString(err)) + "\n"); } -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, @@ -266,7 +251,6 @@ void dcn_v2_cuda_forward(cublasStatus_t stat, cublasHandle_t handle, const int height = in_h; const int width = in_w; - const int channels_out = out_c; 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, ones, k, bias, k, &beta, output, n); - if (stat != CUBLAS_STATUS_SUCCESS) { - printf ("CUBLAS initialization failed\n"); - return ; - } + if (stat != CUBLAS_STATUS_SUCCESS) + FatalError("CUBLAS initialization failed\n"); - modulated_deformable_im2col_cuda(stream, + modulatedDeformableIm2colCuda(stream, input, offset, mask, 1, channels, height, width, height_out, width_out, deformable_group, columns); - // modulated_deformable_im2col_cuda2(stream, + // modulatedDeformableIm2colCudaGeneralVersion(stream, // input, offset, // mask, // 1, channels, height, width, @@ -310,8 +292,7 @@ void dcn_v2_cuda_forward(cublasStatus_t stat, cublasHandle_t handle, columns, n, weight, k, &beta, output, n); - if (stat != CUBLAS_STATUS_SUCCESS) { - printf ("CUBLAS initialization failed\n"); - return ; - } + if (stat != CUBLAS_STATUS_SUCCESS) + FatalError("CUBLAS initialization failed\n"); + }