tensorrt之conv2d计算详解
环境搭建
环境搭建参考文档 TensorRT使用入门
简介
本节主要介绍使用tensorrt来计算conv2d，主要介绍
addConvolutionNd接口的使用，以及IConvolutionLayer中参数的设置。
代码
建议先熟悉环境搭建文档中的demo，熟悉基本的c++ api使用。主要代码如下。
先引入必要的头文件
#include "NvInfer.h"
#include <iostream>
#include <cuda_runtime_api.h>
#include <vector>
#include <sstream>
#include <assert.h>
using namespace nvinfer1;
#define DEFAULT_VALUE 1.0
定义logger
class Logger : public ILogger
{
public:
void log(Severity severity, const char* msg) noexcept override
{
// suppress info-level messages
if (severity <= Severity::kWARNING)
std::cout << msg << std::endl;
}
};
定义一些辅助函数
size_t ProductOfDims(Dims dims) {
size_t result = 1;
for(size_t i = 0; i < dims.nbDims; i++) {
result *= dims.d[i];
}
return result;
}
std::string DimsToStr(Dims dims) {
std::stringstream ss;
for(size_t i = 0; i < dims.nbDims; i++) {
ss << dims.d[i] << " ";
}
return ss.str();
}
Dims4 CalculateConv2dOutput(Dims4 input, Dims4 filter, DimsHW kernel_size, DimsHW stride, DimsHW paddings) {
/******dataformat is NCHW******/
assert(input.d[1] == filter.d[1]); // assert same channel value
assert(filter.d[2] == kernel_size.d[0]);
assert(filter.d[3] == kernel_size.d[1]);
Dims4 output;
output.d[0] = input.d[0]; //batch
output.d[1] = filter.d[0]; //output channel
int32_t kh = kernel_size.d[0];
int32_t kw = kernel_size.d[1];
int32_t sh = stride.d[0];
int32_t sw = stride.d[1];
output.d[2] = ((input.d[2] + 2 * paddings.d[0] - kh) / sh) + 1; //cal output height
output.d[3] = ((input.d[3] + 2 * paddings.d[1] - kw) / sw) + 1; //cal output weight
return output;
}
三个函数的意思显而易见，ProductOfDims计算输入Dims各维度数值的乘积（即元素个数），CalculateConv2dOutput为了计算conv2d的输出shape。注意：CalculateConv2dOutput函数目前只考虑下paddings，未考虑prepadding, postpadding以及padding mode。这在后面会讨论到。
创建一个Network
int main() {
Logger logger;
// Create a Network Definition
IBuilder* builder = createInferBuilder(logger);
uint32_t flag = 1U << static_cast<uint32_t>(NetworkDefinitionCreationFlag::kEXPLICIT_BATCH);
INetworkDefinition* network = builder->createNetworkV2(flag);
定义conv2d的输入以及一些必要属性的shape信息
Dims4 input_shape{3, 3, 8, 8};
Dims4 filter_shape{3, 3, 2, 2};
DimsHW kernel_size{2, 2};
DimsHW stride{1, 1};
DimsHW paddings{1, 1}; //2d padding (padding is symmetric)
Dims4 output_shape = CalculateConv2dOutput(input_shape, filter_shape, kernel_size, stride, paddings);
// Add the Input layer to the network
auto input_data = network->addInput("input", DataType::kFLOAT, input_shape);
添加一个convolution layer，设置stride， padding，
// Add the Convolution layer with hidden layer input nodes, strides and weights for filter and bias.
std::vector<float>filter(ProductOfDims(filter_shape), DEFAULT_VALUE);
Weights filter_w{DataType::kFLOAT, filter.data(), filter.size()};
Weights bias_w{DataType::kFLOAT, nullptr, 0}; // no bias
int32_t output_channel = filter_shape.d[0];
auto conv2d = network->addConvolutionNd(*input_data, output_channel, kernel_size, filter_w, bias_w);
conv2d->setStrideNd(stride);
conv2d->setPaddingNd(paddings);
完整代码见附录1
设置prepadding和postpadding
2d conv的padding有四个方向，上左下右。
前面demo中设置的padding是二维的，即HW两个方向，由于IConvolutionLayer中padding这个属性是对称的，所以padding[0]是上下padding， padding[1]左右padding。
但是实际应用中padding不一定是对称的，即上下左右padding可能各不相同。这是就需要设置IConvolutionLayer中的prepadding和postpadding属性。对于2d conv，prepadding[0]和prepadding[1]分别表示上padding，左padding，postpadding[0]和postpadding[1]分别表示下padding，右padding。
设置prepadding 和 postpadding 的demo如附录2，
我们简化下输入：
Dims4 input_shape{1, 1, 2, 2};
Dims4 filter_shape{1, 1, 2, 2};
DimsHW kernel_size{2, 2};
DimsHW stride{1, 1};
Dims2 prepadding{1,0};
Dims2 postpadding{0,1};
这里定义了上padding和右padding为1。同时conv2d的layer设置对应的属性
auto conv2d = network->addConvolutionNd(*input_data, output_channel, kernel_size, filter_w, bias_w);
conv2d->setStrideNd(stride);
conv2d->setPrePadding(prepadding);
conv2d->setPostPadding(postpadding);
同时注意CalculateConv2dOutput的修改
Dims4 CalculateConv2dOutput(Dims4 input, Dims4 filter, DimsHW kernel_size, DimsHW stride, Dims2 prepadding, Dims2 postpadding) {
/******dataformat is NCHW******/
assert(input.d[1] == filter.d[1]); // assert same channel value
assert(filter.d[2] == kernel_size.d[0]);
assert(filter.d[3] == kernel_size.d[1]);
Dims4 output;
output.d[0] = input.d[0]; //batch
output.d[1] = filter.d[0]; //output channel
int32_t kh = kernel_size.d[0];
int32_t kw = kernel_size.d[1];
int32_t sh = stride.d[0];
int32_t sw = stride.d[1];
output.d[2] = ((input.d[2] + prepadding.d[0] + postpadding.d[0] - kh) / sh) + 1; //cal output height
output.d[3] = ((input.d[3] + prepadding.d[1] + postpadding.d[1] - kw) / sw) + 1; //cal output weight
return output;
}
最后的运行结果如下：
output shape : 1 1 2 2
output data :
2 1 4 2
完整代码见附录2
prepadding、postpadding与paddingNd冲突时
通过前面的介绍可知，IConvolutionLayer有三个接口设置不同情况的padding，即
conv2d->setPaddingNd(padding);
conv2d->setPrePadding(prepadding);
conv2d->setPostPadding(postpadding);
如果三个同时设置了，靠后的设置会覆盖之前的设置。比如上面这种情况，prepadding和postpadding有效，padding无效。读者可自信实验验证。根本原因是：调用setPaddingNd时，内部去调用接口setPrePadding和setPostPadding。也就是说，setPaddingNd接口更改的底层数据和setPrePadding、setPostPadding是同一份数据（prepadding， postpadding）。
我们可以通过一个小case验证一下：
DimsHW padding{1, 1};
conv2d->setPaddingNd(padding);
std::cout << DimsToStr(conv2d->getPaddingNd()) << "\n";
std::cout << DimsToStr(conv2d->getPrePadding()) << "\n";
std::cout << DimsToStr(conv2d->getPostPadding()) << "\n";
输出为：
1 1
1 1
1 1
设置padding mode
IConvolutionLayer设置padding mode的api是
conv2d->setPaddingMode(PaddingMode::kEXPLICIT_ROUND_DOWN); //defalt mode
PaddingMode主要包括三种情况，kEXPLICIT_ROUND_DOWN是默认值，此时计算方式就是前面写的CalculateConv2dOutput中的计算方式。
PaddingMode详细信息可以参考官方文档：
PaddingMode官方文档
注意，从图中公式可知，padding mode为某些场景时，会自动计算prepadding，postpadding(虚拟的值，实际计算o时，如ceil(I/S)，不需要pre/post padding的值)，这种情况设置prepadding， postpadding的优先级更高。（原文 padding mode takes precedence if setPaddingMode() and setPrePadding() are also used.）意思是当你设置的padding mode需要自己推到prepadding， postpadding时，通过setPaddingNd，setPrePadding，setPostPadding这些接口设置的prepadding，postpadding值无效。读者可自行验证。
总结
经过上述的描述后，我们知道，关于padding，本质内部只有两个数据结构，prepadding和postpadding，setPaddingNd(),setPrePadding(),setPostPadding()这三个API都会去设置这两个数据结构，同时，当padding mode为某些场景时，也会去自动设置这两个数据，并且拥有更高的优先值。因为基于上述讨论，我们可以完善下CalculateConv2dOutput()这个函数的逻辑。
Dims4 CalculateConv2dOutput(Dims4 input, Dims4 filter, DimsHW kernel_size, DimsHW stride, Dims2 prepadding,
Dims2 postpadding, PaddingMode paddingmode = PaddingMode::kEXPLICIT_ROUND_DOWN) {
/******dataformat is NCHW******/
assert(input.d[1] == filter.d[1]); // assert same channel value
assert(filter.d[2] == kernel_size.d[0]);
assert(filter.d[3] == kernel_size.d[1]);
Dims4 output;
output.d[0] = input.d[0]; //batch
output.d[1] = filter.d[0]; //output channel
int32_t kh = kernel_size.d[0];
int32_t kw = kernel_size.d[1];
int32_t sh = stride.d[0];
int32_t sw = stride.d[1];
int32_t uppadding = prepadding.d[0];
int32_t leftpadding = prepadding.d[1];
int32_t downpadding = postpadding.d[0];
int32_t rightpadding = postpadding.d[1];
switch (paddingmode)
{
case PaddingMode::kEXPLICIT_ROUND_UP:
output.d[2] = ((input.d[2] + uppadding + downpadding - kh + sh -1 )/ sh) + 1;
output.d[3] = ((input.d[3] + leftpadding + rightpadding - kw + sw -1) / sw) + 1;
break;
case PaddingMode::kSAME_LOWER:
output.d[2] = (input.d[2] / sh) + 1;
output.d[3] = (input.d[3] / sw) + 1;
break;
case PaddingMode::kSAME_UPPER:
output.d[2] = (input.d[2] + sh -1 / sh) + 1;
output.d[3] = (input.d[3] + sw -1 / sw) + 1;
break;
default:
output.d[2] = ((input.d[2] + uppadding + downpadding - kh) / sh) + 1;
output.d[3] = ((input.d[3] + leftpadding + rightpadding - kw) / sw) + 1;
break;
}
return output;
}
代码中如
output.d[2] = (input.d[2] / sh) + 1 对应公式 floor(I / S),
output.d[2] = (input.d[2] + sh -1 / sh) + 1 则对应向上取整的 ceil(I / S)
其他paddingmode的计算方式读者可自行添加。
完整测试代码见附录3
附录
附录1
#include "NvInfer.h"
#include <iostream>
#include <cuda_runtime_api.h>
#include <vector>
#include <sstream>
#include <assert.h>
using namespace nvinfer1;
#define DEFAULT_VALUE 1.0
class Logger : public ILogger
{
public:
void log(Severity severity, const char* msg) noexcept override
{
// suppress info-level messages
if (severity <= Severity::kWARNING)
std::cout << msg << std::endl;
}
};
size_t ProductOfDims(Dims dims) {
size_t result = 1;
for(size_t i = 0; i < dims.nbDims; i++) {
result *= dims.d[i];
}
return result;
}
std::string DimsToStr(Dims dims) {
std::stringstream ss;
for(size_t i = 0; i < dims.nbDims; i++) {
ss << dims.d[i] << " ";
}
return ss.str();
}
Dims4 CalculateConv2dOutput(Dims4 input, Dims4 filter, DimsHW kernel_size, DimsHW stride, DimsHW paddings) {
/******dataformat is NCHW******/
assert(input.d[1] == filter.d[1]); // assert same channel value
assert(filter.d[2] == kernel_size.d[0]);
assert(filter.d[3] == kernel_size.d[1]);
Dims4 output;
output.d[0] = input.d[0]; //batch
output.d[1] = filter.d[0]; //output channel
int32_t kh = kernel_size.d[0];
int32_t kw = kernel_size.d[1];
int32_t sh = stride.d[0];
int32_t sw = stride.d[1];
output.d[2] = ((input.d[2] + 2 * paddings.d[0] - kh) / sh) + 1; //cal output height
output.d[3] = ((input.d[3] + 2 * paddings.d[1] - kw) / sw) + 1; //cal output weight
return output;
}
int main() {
Logger logger;
// Create a Network Definition
IBuilder* builder = createInferBuilder(logger);
uint32_t flag = 1U << static_cast<uint32_t>(NetworkDefinitionCreationFlag::kEXPLICIT_BATCH);
INetworkDefinition* network = builder->createNetworkV2(flag);
Dims4 input_shape{3, 3, 8, 8};
Dims4 filter_shape{3, 3, 2, 2};
DimsHW kernel_size{2, 2};
DimsHW stride{1, 1};
DimsHW paddings{1, 1}; //2d padding (padding is symmetric)
Dims4 output_shape = CalculateConv2dOutput(input_shape, filter_shape, kernel_size, stride, paddings);
// Add the Input layer to the network
auto input_data = network->addInput("input", DataType::kFLOAT, input_shape);
// Add the Convolution layer with hidden layer input nodes, strides and weights for filter and bias.
std::vector<float>filter(ProductOfDims(filter_shape), DEFAULT_VALUE);
Weights filter_w{DataType::kFLOAT, filter.data(), filter.size()};
Weights bias_w{DataType::kFLOAT, nullptr, 0}; // no bias
int32_t output_channel = filter_shape.d[0];
auto conv2d = network->addConvolutionNd(*input_data, output_channel, kernel_size, filter_w, bias_w);
conv2d->setStrideNd(stride);
conv2d->setPaddingNd(paddings);
// Add a name for the output of the conv2d layer so that the tensor can be bound to a memory buffer at inference time:
conv2d->getOutput(0)->setName("output");
// Mark it as the output of the entire network:
network->markOutput(*conv2d->getOutput(0));
// Building an Engine(optimize the network)
IBuilderConfig* config = builder->createBuilderConfig();
IHostMemory* serializedModel = builder->buildSerializedNetwork(*network, *config);
IRuntime* runtime = createInferRuntime(logger);
ICudaEngine* engine = runtime->deserializeCudaEngine(serializedModel->data(), serializedModel->size());
// Prepare input_data
int32_t inputIndex = engine->getBindingIndex("input");
int32_t outputIndex = engine->getBindingIndex("output");
std::vector<float> input(ProductOfDims(input_shape), DEFAULT_VALUE);
std::vector<float> output(ProductOfDims(output_shape));
void *GPU_input_Buffer_ptr; // a host ptr point to a GPU buffer
void *GPU_output_Buffer_ptr; // a host ptr point to a GPU buffer
void* buffers[2];
cudaMalloc(&GPU_input_Buffer_ptr, sizeof(float)*input.size()); //malloc gpu buffer for input
cudaMalloc(&GPU_output_Buffer_ptr, sizeof(float)*output.size()); //malloc gpu buffer for output
cudaMemcpy(GPU_input_Buffer_ptr, input.data(), input.size()*sizeof(float), cudaMemcpyHostToDevice); // copy input data from cpu to gpu
buffers[inputIndex] = static_cast<void*>(GPU_input_Buffer_ptr);
buffers[outputIndex] = static_cast<void*>(GPU_output_Buffer_ptr);
// Performing Inference
IExecutionContext *context = engine->createExecutionContext();
context->executeV2(buffers);
// copy result data from gpu to cpu
cudaMemcpy(output.data(), GPU_output_Buffer_ptr, output.size()*sizeof(float), cudaMemcpyDeviceToHost);
// display output
std::cout << "output shape : " << DimsToStr(output_shape) << "\n";
std::cout << "output data : \n";
for(auto i : output)
std::cout << i << " ";
std::cout << std::endl;
}
附录2
#include "NvInfer.h"
#include <iostream>
#include <cuda_runtime_api.h>
#include <vector>
#include <sstream>
#include <assert.h>
using namespace nvinfer1;
#define DEFAULT_VALUE 1.0
class Logger : public ILogger
{
public:
void log(Severity severity, const char* msg) noexcept override
{
// suppress info-level messages
if (severity <= Severity::kWARNING)
std::cout << msg << std::endl;
}
};
size_t ProductOfDims(Dims dims) {
size_t result = 1;
for(size_t i = 0; i < dims.nbDims; i++) {
result *= dims.d[i];
}
return result;
}
std::string DimsToStr(Dims dims) {
std::stringstream ss;
for(size_t i = 0; i < dims.nbDims; i++) {
ss << dims.d[i] << " ";
}
return ss.str();
}
Dims4 CalculateConv2dOutput(Dims4 input, Dims4 filter, DimsHW kernel_size, DimsHW stride, Dims2 prepadding, Dims2 postpadding) {
/******dataformat is NCHW******/
assert(input.d[1] == filter.d[1]); // assert same channel value
assert(filter.d[2] == kernel_size.d[0]);
assert(filter.d[3] == kernel_size.d[1]);
Dims4 output;
output.d[0] = input.d[0]; //batch
output.d[1] = filter.d[0]; //output channel
int32_t kh = kernel_size.d[0];
int32_t kw = kernel_size.d[1];
int32_t sh = stride.d[0];
int32_t sw = stride.d[1];
output.d[2] = ((input.d[2] + prepadding.d[0] + postpadding.d[0] - kh) / sh) + 1; //cal output height
output.d[3] = ((input.d[3] + prepadding.d[1] + postpadding.d[1] - kw) / sw) + 1; //cal output weight
return output;
}
int main() {
Logger logger;
// Create a Network Definition
IBuilder* builder = createInferBuilder(logger);
uint32_t flag = 1U << static_cast<uint32_t>(NetworkDefinitionCreationFlag::kEXPLICIT_BATCH);
INetworkDefinition* network = builder->createNetworkV2(flag);
Dims4 input_shape{1, 1, 2, 2};
Dims4 filter_shape{1, 1, 2, 2};
DimsHW kernel_size{2, 2};
DimsHW stride{1, 1};
Dims2 prepadding{1,0};
Dims2 postpadding{0,1};
Dims4 output_shape = CalculateConv2dOutput(input_shape, filter_shape, kernel_size, stride, prepadding, postpadding);
// Add the Input layer to the network
auto input_data = network->addInput("input", DataType::kFLOAT, input_shape);
// Add the Convolution layer with hidden layer input nodes, strides and weights for filter and bias.
std::vector<float>filter(ProductOfDims(filter_shape), DEFAULT_VALUE);
Weights filter_w{DataType::kFLOAT, filter.data(), filter.size()};
Weights bias_w{DataType::kFLOAT, nullptr, 0}; // no bias
int32_t output_channel = filter_shape.d[0];
auto conv2d = network->addConvolutionNd(*input_data, output_channel, kernel_size, filter_w, bias_w);
conv2d->setStrideNd(stride);
conv2d->setPrePadding(prepadding);
conv2d->setPostPadding(postpadding);
// Add a name for the output of the conv2d layer so that the tensor can be bound to a memory buffer at inference time:
conv2d->getOutput(0)->setName("output");
// Mark it as the output of the entire network:
network->markOutput(*conv2d->getOutput(0));
// Building an Engine(optimize the network)
IBuilderConfig* config = builder->createBuilderConfig();
IHostMemory* serializedModel = builder->buildSerializedNetwork(*network, *config);
IRuntime* runtime = createInferRuntime(logger);
ICudaEngine* engine = runtime->deserializeCudaEngine(serializedModel->data(), serializedModel->size());
// Prepare input_data
int32_t inputIndex = engine->getBindingIndex("input");
int32_t outputIndex = engine->getBindingIndex("output");
std::vector<float> input(ProductOfDims(input_shape), DEFAULT_VALUE);
std::vector<float> output(ProductOfDims(output_shape));
void *GPU_input_Buffer_ptr; // a host ptr point to a GPU buffer
void *GPU_output_Buffer_ptr; // a host ptr point to a GPU buffer
void* buffers[2];
cudaMalloc(&GPU_input_Buffer_ptr, sizeof(float)*input.size()); //malloc gpu buffer for input
cudaMalloc(&GPU_output_Buffer_ptr, sizeof(float)*output.size()); //malloc gpu buffer for output
cudaMemcpy(GPU_input_Buffer_ptr, input.data(), input.size()*sizeof(float), cudaMemcpyHostToDevice); // copy input data from cpu to gpu
buffers[inputIndex] = static_cast<void*>(GPU_input_Buffer_ptr);
buffers[outputIndex] = static_cast<void*>(GPU_output_Buffer_ptr);
// Performing Inference
IExecutionContext *context = engine->createExecutionContext();
context->executeV2(buffers);
// copy result data from gpu to cpu
cudaMemcpy(output.data(), GPU_output_Buffer_ptr, output.size()*sizeof(float), cudaMemcpyDeviceToHost);
// display output
std::cout << "output shape : " << DimsToStr(output_shape) << "\n";
std::cout << "output data : \n";
for(auto i : output)
std::cout << i << " ";
std::cout << std::endl;
}
附录3
#include "NvInfer.h"
#include <iostream>
#include <cuda_runtime_api.h>
#include <vector>
#include <sstream>
#include <assert.h>
using namespace nvinfer1;
#define DEFAULT_VALUE 1.0
class Logger : public ILogger
{
public:
void log(Severity severity, const char* msg) noexcept override
{
// suppress info-level messages
if (severity <= Severity::kWARNING)
std::cout << msg << std::endl;
}
};
size_t ProductOfDims(Dims dims) {
size_t result = 1;
for(size_t i = 0; i < dims.nbDims; i++) {
result *= dims.d[i];
}
return result;
}
std::string DimsToStr(Dims dims) {
std::stringstream ss;
for(size_t i = 0; i < dims.nbDims; i++) {
ss << dims.d[i] << " ";
}
return ss.str();
}
Dims4 CalculateConv2dOutput(Dims4 input, Dims4 filter, DimsHW kernel_size, DimsHW stride, Dims2 prepadding,
Dims2 postpadding, PaddingMode paddingmode = PaddingMode::kEXPLICIT_ROUND_DOWN) {
/******dataformat is NCHW******/
assert(input.d[1] == filter.d[1]); // assert same channel value
assert(filter.d[2] == kernel_size.d[0]);
assert(filter.d[3] == kernel_size.d[1]);
Dims4 output;
output.d[0] = input.d[0]; //batch
output.d[1] = filter.d[0]; //output channel
int32_t kh = kernel_size.d[0];
int32_t kw = kernel_size.d[1];
int32_t sh = stride.d[0];
int32_t sw = stride.d[1];
int32_t uppadding = prepadding.d[0];
int32_t leftpadding = prepadding.d[1];
int32_t downpadding = postpadding.d[0];
int32_t rightpadding = postpadding.d[1];
switch (paddingmode)
{
case PaddingMode::kEXPLICIT_ROUND_UP:
output.d[2] = ((input.d[2] + uppadding + downpadding - kh + sh -1 )/ sh) + 1;
output.d[3] = ((input.d[3] + leftpadding + rightpadding - kw + sw -1) / sw) + 1;
break;
case PaddingMode::kSAME_LOWER:
output.d[2] = (input.d[2] / sh) + 1;
output.d[3] = (input.d[3] / sw) + 1;
break;
case PaddingMode::kSAME_UPPER:
output.d[2] = (input.d[2] + sh -1 / sh) + 1;
output.d[3] = (input.d[3] + sw -1 / sw) + 1;
break;
default:
output.d[2] = ((input.d[2] + uppadding + downpadding - kh) / sh) + 1;
output.d[3] = ((input.d[3] + leftpadding + rightpadding - kw) / sw) + 1;
break;
}
return output;
}
int main() {
Logger logger;
// Create a Network Definition
IBuilder* builder = createInferBuilder(logger);
uint32_t flag = 1U << static_cast<uint32_t>(NetworkDefinitionCreationFlag::kEXPLICIT_BATCH);
INetworkDefinition* network = builder->createNetworkV2(flag);
Dims4 input_shape{1, 1, 4, 4};
Dims4 filter_shape{1, 1, 2, 2};
DimsHW kernel_size{2, 2};
DimsHW stride{1, 1};
DimsHW padding{1, 1};
Dims2 prepadding{1,0};
Dims2 postpadding{0,1};
Dims4 output_shape = CalculateConv2dOutput(input_shape, filter_shape, kernel_size, stride, prepadding, postpadding);
// Add the Input layer to the network
auto input_data = network->addInput("input", DataType::kFLOAT, input_shape);
// Add the Convolution layer with hidden layer input nodes, strides and weights for filter and bias.
std::vector<float>filter(ProductOfDims(filter_shape), DEFAULT_VALUE);
Weights filter_w{DataType::kFLOAT, filter.data(), filter.size()};
Weights bias_w{DataType::kFLOAT, nullptr, 0}; // no bias
int32_t output_channel = filter_shape.d[0];
auto conv2d = network->addConvolutionNd(*input_data, output_channel, kernel_size, filter_w, bias_w);
conv2d->setStrideNd(stride);
// conv2d->setPaddingNd(padding);
// conv2d->setPrePadding(prepadding);
// conv2d->setPostPadding(postpadding);
conv2d->setPaddingMode(PaddingMode::kSAME_UPPER); //defalt mode
// Add a name for the output of the conv2d layer so that the tensor can be bound to a memory buffer at inference time:
conv2d->getOutput(0)->setName("output");
// Mark it as the output of the entire network:
network->markOutput(*conv2d->getOutput(0));
// Building an Engine(optimize the network)
IBuilderConfig* config = builder->createBuilderConfig();
IHostMemory* serializedModel = builder->buildSerializedNetwork(*network, *config);
IRuntime* runtime = createInferRuntime(logger);
ICudaEngine* engine = runtime->deserializeCudaEngine(serializedModel->data(), serializedModel->size());
// Prepare input_data
int32_t inputIndex = engine->getBindingIndex("input");
int32_t outputIndex = engine->getBindingIndex("output");
std::vector<float> input(ProductOfDims(input_shape), DEFAULT_VALUE);
std::vector<float> output(ProductOfDims(output_shape));
void *GPU_input_Buffer_ptr; // a host ptr point to a GPU buffer
void *GPU_output_Buffer_ptr; // a host ptr point to a GPU buffer
void* buffers[2];
cudaMalloc(&GPU_input_Buffer_ptr, sizeof(float)*input.size()); //malloc gpu buffer for input
cudaMalloc(&GPU_output_Buffer_ptr, sizeof(float)*output.size()); //malloc gpu buffer for output
cudaMemcpy(GPU_input_Buffer_ptr, input.data(), input.size()*sizeof(float), cudaMemcpyHostToDevice); // copy input data from cpu to gpu
buffers[inputIndex] = static_cast<void*>(GPU_input_Buffer_ptr);
buffers[outputIndex] = static_cast<void*>(GPU_output_Buffer_ptr);
// Performing Inference
IExecutionContext *context = engine->createExecutionContext();
context->executeV2(buffers);
// copy result data from gpu to cpu
cudaMemcpy(output.data(), GPU_output_Buffer_ptr, output.size()*sizeof(float), cudaMemcpyDeviceToHost);
// display output
std::cout << "output shape : " << DimsToStr(output_shape) << "\n";
std::cout << "output data : \n";
for(auto i : output)
std::cout << i << " ";
std::cout << std::endl;
}
1832
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
