TensorRT 분석 — ONNX→Engine·FP16·INT8·DLA·Multi-Stream

1
PyTorch/TF model
2
   ↓ export
3
ONNX (graph representation)
4
   ↓ trtexec / Builder API
5
.engine (device-specific binary)
6
   ↓ load at runtime
7
inference

1
#include <NvInfer.h>
2
#include <NvOnnxParser.h>
3

4
class Logger : public nvinfer1::ILogger {
5
    void log(Severity s, const char *msg) noexcept override {
6
        if (s <= Severity::kWARNING) std::cerr << msg << "\n";
7
    }
8
} g_logger;
9

10
void build_engine(const char *onnx_path, const char *engine_path) {
11
    auto builder = nvinfer1::createInferBuilder(g_logger);
12
    auto flag    = 1U << int(nvinfer1::NetworkDefinitionCreationFlag::kEXPLICIT_BATCH);
13
    auto network = builder->createNetworkV2(flag);
14

15
    auto parser = nvonnxparser::createParser(*network, g_logger);
16
    parser->parseFromFile(onnx_path, int(nvinfer1::ILogger::Severity::kWARNING));
17

18
    auto config = builder->createBuilderConfig();
19
    config->setMemoryPoolLimit(nvinfer1::MemoryPoolType::kWORKSPACE, 1ULL << 30);
20
    config->setFlag(nvinfer1::BuilderFlag::kFP16);
21

22
    auto serialized = builder->buildSerializedNetwork(*network, *config);
23

24
    std::ofstream f(engine_path, std::ios::binary);
25
    f.write(static_cast<const char*>(serialized->data()), serialized->size());
26
}

1
auto runtime = nvinfer1::createInferRuntime(g_logger);
2
auto engine  = runtime->deserializeCudaEngine(buf.data(), buf.size());
3
auto ctx     = engine->createExecutionContext();
4

5
void *bufs[2];
6
cudaMalloc(&bufs[0], input_bytes);
7
cudaMalloc(&bufs[1], output_bytes);
8

9
ctx->setTensorAddress("input",  bufs[0]);
10
ctx->setTensorAddress("output", bufs[1]);
11

12
cudaStream_t stream;
13
cudaStreamCreate(&stream);
14

15
cudaMemcpyAsync(bufs[0], host_input,  input_bytes, cudaMemcpyHostToDevice, stream);
16
ctx->enqueueV3(stream);
17
cudaMemcpyAsync(host_output, bufs[1], output_bytes, cudaMemcpyDeviceToHost, stream);
18
cudaStreamSynchronize(stream);

1
class Int8Calibrator : public nvinfer1::IInt8EntropyCalibrator2 {
2
    std::vector<std::string> files_;
3
    int idx_ = 0;
4
    void *device_input_;
5
    size_t input_size_;
6
public:
7
    int32_t getBatchSize() const noexcept override { return 1; }
8
    bool getBatch(void *bindings[], const char *names[],
9
                  int n_bindings) noexcept override {
10
        if (idx_ >= files_.size()) return false;
11
        load_image(files_[idx_++], host_buf_);
12
        cudaMemcpy(device_input_, host_buf_, input_size_,
13
                   cudaMemcpyHostToDevice);
14
        bindings[0] = device_input_;
15
        return true;
16
    }
17
    /* readCalibrationCache / writeCalibrationCache는 파일 IO만 */
18
};
19

20
config->setFlag(nvinfer1::BuilderFlag::kINT8);
21
config->setInt8Calibrator(&calibrator);

1
constexpr int N = 4;
2
cudaStream_t streams[N];
3
nvinfer1::IExecutionContext *ctxs[N];
4
void *bufs[N][2];
5

6
for (int i = 0; i < N; i++) {
7
    cudaStreamCreate(&streams[i]);
8
    ctxs[i] = engine->createExecutionContext();
9
    cudaMalloc(&bufs[i][0], in_bytes);
10
    cudaMalloc(&bufs[i][1], out_bytes);
11
}
12

13
for (int f = 0; f < n_frames; f++) {
14
    int s = f % N;
15
    cudaMemcpyAsync(bufs[s][0], inputs[f], in_bytes,
16
                    cudaMemcpyHostToDevice, streams[s]);
17
    ctxs[s]->setTensorAddress("input",  bufs[s][0]);
18
    ctxs[s]->setTensorAddress("output", bufs[s][1]);
19
    ctxs[s]->enqueueV3(streams[s]);
20
    cudaMemcpyAsync(outputs[f], bufs[s][1], out_bytes,
21
                    cudaMemcpyDeviceToHost, streams[s]);
22
}
23
for (int i = 0; i < N; i++) cudaStreamSynchronize(streams[i]);

1
config->setDefaultDeviceType(nvinfer1::DeviceType::kDLA);
2
config->setDLACore(0);
3
config->setFlag(nvinfer1::BuilderFlag::kGPU_FALLBACK);
4
config->setFlag(nvinfer1::BuilderFlag::kINT8);

1
class MyPlugin : public nvinfer1::IPluginV2DynamicExt {
2
public:
3
    int enqueue(const nvinfer1::PluginTensorDesc *inDesc,
4
                const nvinfer1::PluginTensorDesc *outDesc,
5
                const void *const *ins, void *const *outs,
6
                void *ws, cudaStream_t stream) noexcept override {
7
        const int n = inDesc[0].dims.d[0] * inDesc[0].dims.d[1];
8
        my_kernel<<<(n + 255) / 256, 256, 0, stream>>>(
9
            (const float*)ins[0], (float*)outs[0], n);
10
        return 0;
11
    }
12
    /* getOutputDimensions / supportsFormatCombination / clone 등 구현 */
13
};

1
runtime->deserializeCudaEngine(buf, sz);
2
/* device mismatch → nullptr 또는 runtime error */

1
auto config = builder->createBuilderConfig();
2
auto engine = builder->buildSerializedNetwork(*network, *config);
3
/* FP32 — Tensor core 못 씀 */

1
실내 calibration → 야외 deploy → accuracy 10% ↓

1
opset 17 model → 일부 op 미지원 → fallback or error

1
ctx->enqueueV3(default_stream);   /* GPU 30% */

1
config->setDefaultDeviceType(DeviceType::kDLA);
2
/* op unsupported → build fail */