當(dāng)前位置：首頁 > 编程资源 > 编程问答 >内容正文

编程问答

TensorRT(6)-INT8 inference

發(fā)布時間：2024/9/27 编程问答 42 豆豆

生活随笔收集整理的這篇文章主要介紹了 TensorRT(6)-INT8 inference 小編覺得挺不錯的,現(xiàn)在分享給大家,幫大家做個參考.

這一節(jié)通過官方例程介紹 INT8 inference mode.

例程位于?/usr/src/tensorrt/samples/sampleINT8?，是基于mnist的，大體流程是一致的。

流程同樣是 build(Calibration )->deploy，只不過在build時多了一個校準(zhǔn)的操作。

注意以下幾點：

1 網(wǎng)絡(luò)定義

定義網(wǎng)絡(luò)時，注意這個地方傳進(jìn)去的dataType，如果使用FP16 inference 則傳進(jìn)去的是FP16，也就是kHALF；但如果是使用INT8 inference的話，這個地方傳進(jìn)去的是kFLOAT，也就是 FP32，這是因為INT8 需要先用FP32的精度來確定轉(zhuǎn)換系數(shù)，TensorRT自己會在內(nèi)部轉(zhuǎn)換成INT8。

const IBlobNameToTensor* blobNameToTensor =

parser->parse(locateFile(deployFile).c_str(),

locateFile(modelFile).c_str(),

*network,

DataType::kFLOAT);

這個看起來就跟使用FP32是一樣的流程，INT8 MODE inference的輸入和輸出都是 FP32的。

（After the network has been built, it can be used just like an FP32 network, for example, inputs and outputs remain in 32-bit floating point.）

2 校準(zhǔn)網(wǎng)絡(luò)-Calibrating The Network

校準(zhǔn)網(wǎng)絡(luò)時，比較麻煩的是校準(zhǔn)集的構(gòu)建，作者定義了一個BatchStream class來完成這個操作。BatchStream類有個成員函數(shù)getBatch ()是為了依次讀取 batch file 中的數(shù)據(jù)的。

還有個校準(zhǔn)類 Int8EntropyCalibrator，繼承自 NvInfer.h 中的 IInt8EntropyCalibrator

1	class Int8EntropyCalibrator : public IInt8EntropyCalibrator

這個類里面也有個 getBatch () 成員函數(shù)，實際上調(diào)用的是 BatchStream類的getBatch () ，然后將數(shù)據(jù)從內(nèi)存搬到了顯存，如下：

bool getBatch(void* bindings[], const char* names[], int nbBindings) override

{

if (!mStream.next())

return false;

CHECK(cudaMemcpy(mDeviceInput, mStream.getBatch(), mInputCount * sizeof(float), cudaMemcpyHostToDevice));

assert(!strcmp(names[0], INPUT_BLOB_NAME));

bindings[0] = mDeviceInput;

return true;

}

這個getBatch () 成員函數(shù)在校準(zhǔn)時會被反復(fù)調(diào)用。

生成校準(zhǔn)集時，校準(zhǔn)集的樣本應(yīng)該是已經(jīng)進(jìn)行過一系列預(yù)處理的圖片而不是原始圖片。

校準(zhǔn)類 Int8EntropyCalibrator 和 BatchStream 類的實現(xiàn)說起來比較麻煩，在后面源碼解讀部分直接結(jié)合注釋看源碼吧。

3 builder的配置-Configuring The Builder

只需要在原來builder的基礎(chǔ)上添加以下：

1 2	builder->setInt8Mode(true); builder->setInt8Calibrator(calibrator);

4 batch file的生成-Batch Files For Calibration

例程使用的batch file 已經(jīng)制作好了，位于<TensorRT>/data/mnist/batches?這是一系列二進(jìn)制文件，每個文件包含了 N 個圖片樣本，格式如下：

首先是4個32 bit的整形值，代表 {N, C, H, W},batchsize和圖片dims
然后是N個 {C, H, W}維度的浮點數(shù)據(jù)，代表N個樣本

batch file二進(jìn)制文件的生成有兩種方式：

4.1 使用caffe生成

主要對于使用caffe的用戶，這里干脆直接將官方文檔上的說明拷貝過來好了，比較簡單：

Navigate to the samples data directory and create an INT8 mnist directory:

> cd <TensorRT>/samples/data

> mkdir -p int8/mnist

> cd int8/mnist

Note: If Caffe is not installed anywhere, ensure you clone, checkout, patch, and build Caffe at the specific commit:

> git clone https://github.com/BVLC/caffe.git

> cd caffe

> git checkout 473f143f9422e7fc66e9590da6b2a1bb88e50b2f

> patch -p1 < <TensorRT>/samples/mnist/int8_caffe.patch

> mkdir build

> pushd build

> cmake -DUSE_OPENCV=FALSE -DUSE_CUDNN=OFF ../

> make -j4

> popd

Download the mnist dataset from Caffe and create a link to it:

> bash data/mnist/get_mnist.sh

> bash examples/mnist/create_mnist.sh

> cd ..

> ln -s caffe/examples .

Set the directory to store the batch data, execute Caffe, and link the mnist files:

> mkdir batches

> export TENSORRT_INT8_BATCH_DIRECTORY=batches

> caffe/build/tools/caffe test -gpu 0 -iterations 1000 -model examples/mnist/lenet_train_test.prototxt -weights

> <TensorRT>/samples/mnist/mnist.caffemodel

> ln -s <TensorRT>/samples/mnist/mnist.caffemodel .

> ln -s <TensorRT>/samples/mnist/mnist.prototxt .

Execute sampleINT8 from the bin directory after being built with the following command:

1 2	> ./sample_int8 mnist >

4.2 其他方式生成

對于不用caffe或者模型難以轉(zhuǎn)換成caffemode的用戶，首先要進(jìn)行一系列預(yù)處理，然后按照前面提到的batch file格式生成二進(jìn)制batch file文件，但這個生成過程要自己寫了，不過寫的話應(yīng)該也比較簡單，可以參考caffe中的patch文件中的核心部分：

#define LOG_BATCHES_FOR_INT8_TESTING 1

#if LOG_BATCHES_FOR_INT8_TESTING

static int sBatchId = 0;

char* batch_dump_dir = getenv("TENSORRT_INT8_BATCH_DIRECTORY");

if(batch_dump_dir != 0)

{

char buffer[1000];

sprintf(buffer, "batches/batch%d", sBatchId++);

FILE* file = fopen(buffer, "w");

if(file==0)

abort();

int s[4] = { top_shape[0], top_shape[1], top_shape[2], top_shape[3] };

fwrite(s, sizeof(int), 4, file);

fwrite(top_data, sizeof(float), top_shape[0]*top_shape[1]*top_shape[2]*top_shape[3], file);

fwrite(&top_label[0], sizeof(int), top_shape[0], file);

fclose(file);

}

+#endif

添加上數(shù)據(jù)集的讀取，劃分和預(yù)處理就可以了。

5 校準(zhǔn)算法

從INT8的例程來看，TensorRT 支持兩種方式的校準(zhǔn)，一種就是上節(jié)我們講過的使用相對熵的方式，還有一種是廢棄的校準(zhǔn)算法，校準(zhǔn)時需要設(shè)置兩個參數(shù) cutoff 和 quantile，以下是在GTC2017 上對INT8校準(zhǔn)原理進(jìn)行講解的 Szymon Migacz 對廢棄的校準(zhǔn)算法的解讀：

https://devtalk.nvidia.com/default/topic/1015108/cutoff-and-quantile-parameters-in-tensorrt/

Parameters cutoff and quantile have to be specified only for “l(fā)egacy” calibrator. It’s difficult to set values of cutoff and quantile without running experiments. Our recommended way was to run 2D grid search and look for optimal combination of (cutoff, quantile) for a given network on a given dataset. This was implemented in sampleINT8 shipped with TensorRT 2.0 EA.

New entropy calibrator doesn’t require any external hyperparameters, and it determines quantization thresholds automatically based on the distributions of activations on calibration dataset. In my presentation at GTC I was talking only about the new entropy calibrator, it’s available in TensorRT 2.1 GA.

Szymon Migacz并沒有充分的解釋這兩個參數(shù)，而是說這是 “l(fā)egacy” calibrator中才會用到的參數(shù)，而且在沒有做充分的試驗的情況下，是很難合理地設(shè)置這兩個參數(shù)的。他推薦的做法是針對特定的網(wǎng)絡(luò)結(jié)構(gòu)和數(shù)據(jù)集使用 2D 網(wǎng)格搜索來確定這兩個參數(shù)的取值。而 entropy calibrator ，就是使用相對熵的校準(zhǔn)方法，不需要任何超參數(shù)，而且能夠根據(jù)校準(zhǔn)集上的激活值分布自動確定量化閾值。NVIDIA官方也推薦使用使用相對熵校準(zhǔn)的方式。所以 “l(fā)egacy” calibrator 就不深入研究了。

6 源碼解讀

sampleINT8.cpp:

100

101

102

103

104

105

106

107

108

109

110

111

112

113

114

115

116

117

118

119

120

121

122

123

124

125

126

127

128

129

130

131

132

133

134

135

136

137

138

139

140

141

142

143

144

145

146

147

148

149

150

151

152

153

154

155

156

157

158

159

160

161

162

163

164

165

166

167

168

169

170

171

172

173

174

175

176

177

178

179

180

181

182

183

184

185

186

187

188

189

190

191

192

193

194

195

196

197

198

199

200

201

202

203

204

205

206

207

208

209

210

211

212

213

214

215

216

217

218

219

220

221

222

223

224

225

226

227

228

229

230

231

232

233

234

235

236

237

238

239

240

241

242

243

244

245

246

247

248

249

250

251

252

253

254

255

256

257

258

259

260

261

262

263

264

265

266

267

268

269

270

271

272

273

274

275

276

277

278

279

280

281

282

283

284

285

286

287

288

289

290

291

292

293

294

295

296

297

298

299

300

301

302

303

304

305

306

307

308

309

310

311

312

313

314

315

316

317

318

319

320

321

322

323

324

325

326

327

328

329

330

331

332

333

334

335

336

337

338

339

340

341

342

343

344

345

346

347

348

349

350

351

352

353

354

355

356

357

358

359

360

361

362

363

364

365

366

367

368

369

370

371

372

373

374

375

376

377

378

379

380

381

382

383

384

385

386

387

388

389

390

391

392

393

394

395

396

397

398

399

400

401

402

403

404

405

406

407

408

409

410

411

412

413

414

415

416

417

418

419

420

421

422

423

424

425

426

427

428

429

430

431

432

433

434

435

436

437

438

439

440

441

442

443

444

445

446

447

448

449

450

451

452

453

454

455

456

457

458

459

460

461

462

463

464

465

466

467

468

469

470

471

472

473

474

475

#include <assert.h>

#include <fstream>

#include <sstream>

#include <iostream>

#include <cmath>

#include <sys/stat.h>

#include <cmath>

#include <time.h>

#include <cuda_runtime_api.h>

#include <unordered_map>

#include <algorithm>

#include <float.h>

#include <string.h>

#include <chrono>

#include <iterator>

#include "NvInfer.h"

#include "NvCaffeParser.h"

#include "common.h"

#include "BatchStream.h"

#include "LegacyCalibrator.h"

using namespace nvinfer1;

using namespace nvcaffeparser1;

static Logger gLogger;

// stuff we know about the network and the caffe input/output blobs

const char* INPUT_BLOB_NAME = "data";

const char* OUTPUT_BLOB_NAME = "prob";

const char* gNetworkName{nullptr};

std::string locateFile(const std::string& input)

{

std::vector<std::string> dirs;

dirs.push_back(std::string("data/int8/") + gNetworkName + std::string("/"));

dirs.push_back(std::string("data/") + gNetworkName + std::string("/"));

return locateFile(input, dirs);

}

bool caffeToTRTModel(const std::string& deployFile, // name for caffe prototxt

const std::string& modelFile, // name for model

const std::vector<std::string>& outputs, // network outputs

unsigned int maxBatchSize, // batch size - NB must be at least as large as the batch we want to run with)

DataType dataType,

IInt8Calibrator* calibrator,

nvinfer1::IHostMemory *&trtModelStream)

{

//創(chuàng)建一個builder，傳入自己實現(xiàn)的 gLogger 對象，為了打印信息用

// create the builder

IBuilder* builder = createInferBuilder(gLogger);

//創(chuàng)建一個 network 對象，并創(chuàng)建一個 ICaffeParser 對象，這個對象是用來進(jìn)行模型轉(zhuǎn)換的；此時的 network 對象里面還是空的

// parse the caffe model to populate the network, then set the outputs

INetworkDefinition* network = builder->createNetwork();

ICaffeParser* parser = createCaffeParser();

//判斷當(dāng)前的硬件平臺是否支持 INT8 精度和 FP16 精度，兩者都不支持的話，直接返回 false

if((dataType == DataType::kINT8 && !builder->platformHasFastInt8()) || (dataType == DataType::kHALF && !builder->platformHasFastFp16()))

return false;

// caffemodel到tensorrt的轉(zhuǎn)換, 注意這個地方傳進(jìn)去的dataType，

// 如果使用FP16 inference 則傳進(jìn)去的是FP16，也就是kHALF

// 如果是使用INT8 inference的話，這個地方傳進(jìn)去的是kFLOAT也就是 FP32,

// 因為INT8 需要先用FP32的精度來確定轉(zhuǎn)換系數(shù)，TensorRT自己會在內(nèi)部轉(zhuǎn)換成INT8

const IBlobNameToTensor* blobNameToTensor = parser->parse(locateFile(deployFile).c_str(),

locateFile(modelFile).c_str(),

*network,

dataType == DataType::kINT8 ? DataType::kFLOAT : dataType);

//標(biāo)志輸出tensor

// specify which tensors are outputs

for (auto& s : outputs)

network->markOutput(*blobNameToTensor->find(s.c_str()));

// Build the engine

// 設(shè)置最大 batchsize和工作空間大小 2^30 ,這里是1G

builder->setMaxBatchSize(maxBatchSize);

builder->setMaxWorkspaceSize(1 << 30);

// 設(shè)置平均迭代次數(shù)和最小迭代次數(shù)，這是測量每一層時間的一種策略，即多次迭代求平均值，不過這里只迭代一次

builder->setAverageFindIterations(1);

builder->setMinFindIterations(1);

//同步調(diào)試

builder->setDebugSync(true);

//INT8 MODE or/and FP16 MODE

builder->setInt8Mode(dataType == DataType::kINT8);

builder->setFp16Mode(dataType == DataType::kHALF);

//設(shè)置INT8校準(zhǔn)接口

builder->setInt8Calibrator(calibrator);

// 創(chuàng)建engine

ICudaEngine* engine = builder->buildCudaEngine(*network);

assert(engine);

//銷毀無用對象

// we don't need the network any more, and we can destroy the parser

network->destroy();

parser->destroy();

//序列化到磁盤上，這里實際上是在內(nèi)存中，沒有保存到磁盤

// serialize the engine, then close everything down

trtModelStream = engine->serialize();

engine->destroy();

builder->destroy();

return true;

}

float doInference(IExecutionContext& context, float* input, float* output, int batchSize)

{

//從context恢復(fù)engine

const ICudaEngine& engine = context.getEngine();

//創(chuàng)建engine的時候，會把輸入blob和輸出blob指針放進(jìn)去，engine.getNbBindings() 就是為了獲取輸入和輸出的blob數(shù)目，以便于做檢查

//比如這里，就只有一個輸入和一個輸出，所以檢查時可以這樣檢查 assert(engine.getNbBindings() == 2);

// input and output buffer pointers that we pass to the engine - the engine requires exactly IEngine::getNbBindings(),

// of these, but in this case we know that there is exactly one input and one output.

assert(engine.getNbBindings() == 2);

//每個輸入和輸出blob都需要申請顯存，故：void* buffers[engine.getNbBindings()];

void* buffers[2];

float ms{ 0.0f };

//為了將 buffer中的成員(指針或者地址)分別與輸入/輸出的blob相關(guān)聯(lián)，需要分別獲取輸入輸出blob在engine中的索引

// In order to bind the buffers, we need to know the names of the input and output tensors.

// note that indices are guaranteed to be less than IEngine::getNbBindings()

int inputIndex = engine.getBindingIndex(INPUT_BLOB_NAME),

outputIndex = engine.getBindingIndex(OUTPUT_BLOB_NAME);

//計算輸入輸出shape

// create GPU buffers and a stream

Dims3 inputDims = static_cast<Dims3&&>(context.getEngine().getBindingDimensions(context.getEngine().getBindingIndex(INPUT_BLOB_NAME)));

Dims3 outputDims = static_cast<Dims3&&>(context.getEngine().getBindingDimensions(context.getEngine().getBindingIndex(OUTPUT_BLOB_NAME)));

//計算實際的輸入輸出大小,申請顯存

size_t inputSize = batchSize*inputDims.d[0]*inputDims.d[1]*inputDims.d[2] * sizeof(float), outputSize = batchSize *

outputDims.d[0] * outputDims.d[1] * outputDims.d[2] * sizeof(float);

CHECK(cudaMalloc(&buffers[inputIndex], inputSize));

CHECK(cudaMalloc(&buffers[outputIndex], outputSize));

//從Host (CPU) 拷貝輸入數(shù)據(jù)到 Device(GPU)，也就是從內(nèi)存到顯存

CHECK(cudaMemcpy(buffers[inputIndex], input, inputSize, cudaMemcpyHostToDevice));

//創(chuàng)建一個 cuda 異步流

cudaStream_t stream;

CHECK(cudaStreamCreate(&stream));

//創(chuàng)建一個cuda事件

cudaEvent_t start, end;

CHECK(cudaEventCreateWithFlags(&start, cudaEventBlockingSync));

CHECK(cudaEventCreateWithFlags(&end, cudaEventBlockingSync));

//標(biāo)記stream流，start

cudaEventRecord(start, stream);

//異步執(zhí)行inference，//標(biāo)記stream流，end

context.enqueue(batchSize, buffers, stream, nullptr);

cudaEventRecord(end, stream);

//事件同步

cudaEventSynchronize(end);

//計算start事件和end事件之間的運行時間

cudaEventElapsedTime(&ms, start, end);

//銷毀事件

cudaEventDestroy(start);

cudaEventDestroy(end);

//從Device(GPU) 拷貝輸出數(shù)據(jù)到 Host (CPU)，也就是從顯存到內(nèi)存

CHECK(cudaMemcpy(output, buffers[outputIndex], outputSize, cudaMemcpyDeviceToHost));

//釋放顯存

CHECK(cudaFree(buffers[inputIndex]));

CHECK(cudaFree(buffers[outputIndex]));

//銷毀流對象

CHECK(cudaStreamDestroy(stream));

//返回inference時間

return ms;

}

//計算一個batch 中 top-1或top-5的正確的圖片數(shù)量

//對于輸出來說，一張圖片的輸出對應(yīng)一個 outputSize 維的向量（比如mnist是10維的）

//然而對于標(biāo)簽來說一張圖片的標(biāo)簽是一個0-9之間的數(shù)字

//batchProb是一個batch中的標(biāo)簽向量按順序疊加到一個vector中的，10個數(shù)字一組對應(yīng)一張圖片

//label就這這個batch的標(biāo)簽向量，一個數(shù)字對應(yīng)一張圖片

//outputsize是輸出維度（比如mnist的outputsize=10）

//threshold：兩個取值：1，對應(yīng)top-1；5對應(yīng)top-5

int calculateScore(float* batchProb, float* labels, int batchSize, int outputSize, int threshold)

{

int success = 0;

for (int i = 0; i < batchSize; i++)

{

//獲取每個batch的地址，并獲取預(yù)測向量中與標(biāo)簽相同位置上的真實概率

//舉個例子：假設(shè)threshold=1

//i=0時，prob[0]-prob[9]是batch中的第一張圖片的預(yù)測輸出向量，

//假設(shè)prob[0]-prob[9]的值為{0.1, 0.5, 0.05, 0.05, 0.05, 0.05, 0.05, 0.05, 0.05, 0.05}，這張圖片的label是1.

//那么correct = prob[(int)labels[i]]=prob[1]=0.5，之后判斷的是這個correct是否在top-1或者top-5范圍內(nèi)

//做法是：統(tǒng)計 prob[0]-prob[9]之間比correct更大的值的個數(shù) better，因為如果比correct大的話，最終輸出的肯定是錯的預(yù)測結(jié)果；

//但是由于top-1，top-5允許你出錯的次數(shù)分別為1次和5次，所以只要 better < threshold，就認(rèn)為預(yù)測準(zhǔn)確，success++；

//最后返回success，代表這個batch中按照 top-1 或 top-5的精度來算，預(yù)測對了幾張圖片。

float* prob = batchProb + outputSize*i, correct = prob[(int)labels[i]];

int better = 0;

for (int j = 0; j < outputSize; j++)

if (prob[j] >= correct)

better++;

if (better <= threshold)

success++;

}

return success;

}

class Int8EntropyCalibrator : public IInt8EntropyCalibrator

{

public:

Int8EntropyCalibrator(BatchStream& stream, int firstBatch, bool readCache = true)

: mStream(stream), mReadCache(readCache)

{

DimsNCHW dims = mStream.getDims();

mInputCount = mStream.getBatchSize() * dims.c() * dims.h() * dims.w();

//為 mDeviceInput 申請顯存，跳過前面 firstBatch 個batch

CHECK(cudaMalloc(&mDeviceInput, mInputCount * sizeof(float)));

mStream.reset(firstBatch);

}

/**

* 析構(gòu)函數(shù)，釋放顯存

virtual ~Int8EntropyCalibrator()

{

CHECK(cudaFree(mDeviceInput));

}

int getBatchSize() const override { return mStream.getBatchSize(); }

bool getBatch(void* bindings[], const char* names[], int nbBindings) override

{

if (!mStream.next())

return false;

//將mStream.getBatch()獲取到的數(shù)據(jù)拷貝到 mDeviceInput 中，也就是從內(nèi)存到顯存

CHECK(cudaMemcpy(mDeviceInput, mStream.getBatch(), mInputCount * sizeof(float), cudaMemcpyHostToDevice));

assert(!strcmp(names[0], INPUT_BLOB_NAME));

bindings[0] = mDeviceInput;

return true;

}

/**

* 從文件中讀取校準(zhǔn)數(shù)據(jù)，返回校準(zhǔn)表緩存地址

* @param length 讀取長度

const void* readCalibrationCache(size_t& length) override

{

//首先清空mCalibrationCache

mCalibrationCache.clear();

//從文件中讀取內(nèi)容并放到 mCalibrationCache vector中

std::ifstream input(calibrationTableName(), std::ios::binary);

input >> std::noskipws;

if (mReadCache && input.good())

std::copy(std::istream_iterator<char>(input), std::istream_iterator<char>(), std::back_inserter(mCalibrationCache));

//返回 mCalibrationCache 地址或空指針

length = mCalibrationCache.size();

return length ? &mCalibrationCache[0] : nullptr;

}

/**

* 將校準(zhǔn)數(shù)據(jù)存儲到文件中

* @param cache 校準(zhǔn)數(shù)據(jù)內(nèi)存地址

* @param length 數(shù)據(jù)長度

void writeCalibrationCache(const void* cache, size_t length) override

{

std::ofstream output(calibrationTableName(), std::ios::binary);

output.write(reinterpret_cast<const char*>(cache), length);

}

private:

/**

* 存儲校準(zhǔn)數(shù)據(jù)的文件

* @return 文件名稱

static std::string calibrationTableName()

{

assert(gNetworkName);

return std::string("CalibrationTable") + gNetworkName;

}

//batch流

BatchStream mStream;

//是否從文件中讀取校準(zhǔn)數(shù)據(jù)

bool mReadCache{ true };

//校準(zhǔn)時 GPU接受的數(shù)據(jù)量mInputCount 和數(shù)據(jù)內(nèi)容 mDeviceInput

size_t mInputCount;

void* mDeviceInput{ nullptr };

//存放從文件中讀取到的校準(zhǔn)數(shù)據(jù)，也就是scale_factor 縮放系數(shù)

std::vector<char> mCalibrationCache;

};

/**

* 用于模型評分，包含了caffe模型向ensorRT的轉(zhuǎn)化以及inference的執(zhí)行

* @param batchSize 批尺寸

* @param firstBatch 跳過初始的一些batch

* @param nbScoreBatches 測試的 batch總數(shù)

* @param datatype 以何種精度inference

* @param calibrator 校準(zhǔn)接口

* @param quiet 是否輸出調(diào)試信息

std::pair<float, float> scoreModel(int batchSize, int firstBatch, int nbScoreBatches, DataType datatype, IInt8Calibrator* calibrator, bool quiet = false)

{

IHostMemory *trtModelStream{ nullptr };

// 調(diào)用 caffeToTRTModel 將caffe模型解析為TensorRT

bool valid = false;

if (gNetworkName == std::string("mnist"))

valid = caffeToTRTModel("deploy.prototxt", "mnist_lenet.caffemodel", std::vector < std::string > { OUTPUT_BLOB_NAME }, batchSize, datatype, calibrator, trtModelStream);

else

valid = caffeToTRTModel("deploy.prototxt", std::string(gNetworkName) + ".caffemodel", std::vector < std::string > { OUTPUT_BLOB_NAME }, batchSize, datatype, calibrator, trtModelStream);

// 如果GPU不支持某種精度類型，比如FP16/INT8，則返回（0,0）

if(!valid)

{

std::cout << "Engine could not be created at this precision" << std::endl;

return std::pair<float, float>(0,0);

}

assert(trtModelStream != nullptr);

// 恢復(fù)創(chuàng)建engine，創(chuàng)建上下文環(huán)境

// Create engine and deserialize model.

IRuntime* infer = createInferRuntime(gLogger);

assert(infer != nullptr);

ICudaEngine* engine = infer->deserializeCudaEngine(trtModelStream->data(), trtModelStream->size(), nullptr);

assert(engine != nullptr);

trtModelStream->destroy();

IExecutionContext* context = engine->createExecutionContext();

assert(context != nullptr);

//創(chuàng)建 batch 流對象，并跳過開始的一些batch，共firstBatch個，此處等于100

BatchStream stream(batchSize, nbScoreBatches);

stream.skip(firstBatch);

// output tensor 維度

Dims3 outputDims = static_cast<Dims3&&>(context->getEngine().getBindingDimensions(context->getEngine().getBindingIndex(OUTPUT_BLOB_NAME)));

//確定輸出 tensor 數(shù)據(jù)量大小

int outputSize = outputDims.d[0]*outputDims.d[1]*outputDims.d[2];

int top1{ 0 }, top5{ 0 };

float totalTime{ 0.0f };

//每張圖片都有一個 outputSize 大小的向量(比如 mnist 分類大小為10)，那么一個batch的輸出應(yīng)該為 batchSize * outputSize

std::vector<float> prob(batchSize * outputSize, 0);

//依次對不同的batch進(jìn)行inference，stream.next()獲取下一個batch

while (stream.next())

{

//輸入數(shù)據(jù)：stream.getBatch()，輸出數(shù)據(jù)：prob 每循環(huán)一次就對一個batch的數(shù)據(jù)進(jìn)行測試，這個batch的輸出放在 prob 中

totalTime += doInference(*context, stream.getBatch(), &prob[0], batchSize);

//對每個batch，按照top-1和top-5精度來計算準(zhǔn)確率

top1 += calculateScore(&prob[0], stream.getLabels(), batchSize, outputSize, 1);

top5 += calculateScore(&prob[0], stream.getLabels(), batchSize, outputSize, 5);

//讀取10個batch輸出一個點，讀取800個輸出一個換行符

std::cout << (!quiet && stream.getBatchesRead() % 10 == 0 ? "." : "") << (!quiet && stream.getBatchesRead() % 800 == 0 ? "\n" : "") << std::flush;

}

//統(tǒng)計總共讀到了多少張圖片，并計算top-1和top-5正確率

int imagesRead = stream.getBatchesRead()*batchSize;

float t1 = float(top1) / float(imagesRead), t5 = float(top5) / float(imagesRead);

// 精度和時間，結(jié)果輸出

if (!quiet)

{

std::cout << "\nTop1: " << t1 << ", Top5: " << t5 << std::endl;

std::cout << "Processing " << imagesRead << " images averaged " << totalTime / imagesRead << " ms/image and " << totalTime / stream.getBatchesRead() << " ms/batch." << std::endl;

}

//銷毀無用對象，返回準(zhǔn)確率

context->destroy();

engine->destroy();

infer->destroy();

return std::make_pair(t1, t5);

}

int main(int argc, char** argv)

{

if (argc < 2)

{

std::cout << "Please provide the network as the first argument." << std::endl;

exit(0);

}

gNetworkName = argv[1];

//前 firstScoreBatch 個 batch是用來作為校準(zhǔn)集的，因此在測試時這些是不進(jìn)行測試的

int batchSize = 100, firstScoreBatch = 100, nbScoreBatches = 400; // by default we score over 40K images starting at 10000, so we don't score those used to search calibration

//search變量是LEGACY_CALIBRATION校準(zhǔn)算法中使用的變量，具體作用要看 LegacyCalibrator.h 源碼，因為這個校準(zhǔn)算法nvidia已經(jīng)不推薦使用了，所以這里不深究了

bool search = false;

//校準(zhǔn)算法選擇參考 Nvinfer.h 文件，kENTROPY_CALIBRATION：使用信息熵進(jìn)行校準(zhǔn)；kLEGACY_CALIBRATION，使用以前遺留下來的算法進(jìn)行校準(zhǔn)

// enum class CalibrationAlgoType : int

// {

// kLEGACY_CALIBRATION = 0,

// kENTROPY_CALIBRATION = 1

// };

CalibrationAlgoType calibrationAlgo = CalibrationAlgoType::kENTROPY_CALIBRATION;

// 處理命令行參數(shù)

for (int i = 2; i < argc; i++)

{

if (!strncmp(argv[i], "batch=", 6))

batchSize = atoi(argv[i] + 6);

else if (!strncmp(argv[i], "start=", 6))

firstScoreBatch = atoi(argv[i] + 6);

else if (!strncmp(argv[i], "score=", 6))

nbScoreBatches = atoi(argv[i] + 6);

else if (!strncmp(argv[i], "search", 6))

search = true;

else if (!strncmp(argv[i], "legacy", 6))

calibrationAlgo = CalibrationAlgoType::kLEGACY_CALIBRATION;

else

{

std::cout << "Unrecognized argument " << argv[i] << std::endl;

exit(0);

}

if (calibrationAlgo == CalibrationAlgoType::kENTROPY_CALIBRATION)

{

search = false;

}

//batchsize不能大于128，這是為何？

if (batchSize > 128)

{

std::cout << "Please provide batch size <= 128" << std::endl;

exit(0);

}

//感覺這里寫錯了，應(yīng)該是 50000

if ((firstScoreBatch + nbScoreBatches)*batchSize > 500000)

{

std::cout << "Only 50000 images available" << std::endl;

exit(0);

}

//設(shè)置標(biāo)準(zhǔn)輸出流輸出的精度

std::cout.precision(6);

//用于構(gòu)建校準(zhǔn)集的batch流

//CAL_BATCH_SIZE = 50;NB_CAL_BATCHES = 10; 定義在 LegacyCalibrator.h文件中, 既然廢棄了 LegacyCalibrator，為什么不把常量定義在本文件中

BatchStream calibrationStream(CAL_BATCH_SIZE, NB_CAL_BATCHES);

//FP32精度不需要校準(zhǔn)集，因此最后一個參數(shù)傳入 nullptr

std::cout << "\nFP32 run:" << nbScoreBatches << " batches of size " << batchSize << " starting at " << firstScoreBatch << std::endl;

scoreModel(batchSize, firstScoreBatch, nbScoreBatches, DataType::kFLOAT, nullptr);

//FP16精度不需要校準(zhǔn)集，因此最后一個參數(shù)傳入 nullptr

std::cout << "\nFP16 run:" << nbScoreBatches << " batches of size " << batchSize << " starting at " << firstScoreBatch << std::endl;

scoreModel(batchSize, firstScoreBatch, nbScoreBatches, DataType::kHALF, nullptr);

std::cout << "\nINT8 run:" << nbScoreBatches << " batches of size " << batchSize << " starting at " << firstScoreBatch << std::endl;

if (calibrationAlgo == CalibrationAlgoType::kENTROPY_CALIBRATION)

{

//先構(gòu)建校準(zhǔn)集，然后調(diào)用scoreModel進(jìn)行模型評估，創(chuàng)建engine時傳入了Int8EntropyCalibrator對象calibrator

//FIRST_CAL_SCORE_BATCH = 100; 定義在 LegacyCalibrator.h文件中

Int8EntropyCalibrator calibrator(calibrationStream, FIRST_CAL_BATCH);

scoreModel(batchSize, firstScoreBatch, nbScoreBatches, DataType::kINT8, &calibrator);

}

else

{

//被廢棄的校準(zhǔn)算法，不解釋了

std::pair<double, double> parameters = getQuantileAndCutoff(gNetworkName, search);

Int8LegacyCalibrator calibrator(calibrationStream, FIRST_CAL_BATCH, parameters.first, parameters.second);

scoreModel(batchSize, firstScoreBatch, nbScoreBatches, DataType::kINT8, &calibrator);

}

shutdownProtobufLibrary();

return 0;

}

BatchStream.h，這個源碼看起來還是稍微有點費勁的，還是我C++功底不夠啊，得補(bǔ)。。。

100

101

102

103

104

105

106

107

108

109

110

111

112

113

114

115

116

117

118

119

120

121

122

123

124

125

126

127

128

129

130

131

132

133

134

135

136

137

138

139

140

141

142

143

144

145

146

147

148

149

150

151

152

153

154

155

156

157

158

159

160

161

162

163

164

165

166

167

168

169

170

171

172

173

#ifndef BATCH_STREAM_H

#define BATCH_STREAM_H

#include <vector>

#include <assert.h>

#include <algorithm>

#include "NvInfer.h"

std::string locateFile(const std::string& input);

class BatchStream

{

public:

//構(gòu)造函數(shù)，使用 batchSize 和 maxBatches 初始化 BatchStream 中的 mBatchSize(批尺寸) 和 mMaxBatches(批數(shù)量)

BatchStream(int batchSize, int maxBatches) : mBatchSize(batchSize), mMaxBatches(maxBatches)

{

//讀取第一個batch文件的shape，用于一系列初始化操作

FILE* file = fopen(locateFile(std::string("batches/batch0")).c_str(), "rb");

int d[4];

fread(d, sizeof(int), 4, file);

mDims = nvinfer1::DimsNCHW{ d[0], d[1], d[2], d[3] };

fclose(file);

//單張圖片的大小（總的像素個數(shù)）

mImageSize = mDims.c()*mDims.h()*mDims.w();

//根據(jù)batch文件中的單張圖片大小mImageSize初始化 BatchStream 中的 mBatch 的內(nèi)存空間，初值為0；同理根據(jù)mBatchSize初始化mLabels

//mBatch指的是BatchStream中的batch，batch的個數(shù)為mBatchSize，所以數(shù)據(jù)量總數(shù)為mBatchSize*mImageSize，

//mLabels是BatchStream中的label，總數(shù)就是 mBatchSize

mBatch.resize(mBatchSize*mImageSize, 0);

mLabels.resize(mBatchSize, 0);

//有兩塊專門的內(nèi)存區(qū)域用于存儲讀取到的batch{i} 文件內(nèi)容，就是下面兩個。這兩塊內(nèi)存區(qū)域里的內(nèi)容在后面會被復(fù)制到 mBatch和mLabels中

//mFileBatch指的是讀取到的 batch{i} 文件中的batch，因此總數(shù)為mDims.n()*mDims.c()*mDims.h()*mDims.w()=mDims.n()*mImageSize

//mFileLabels指的是讀取到的 batch{i} 文件中的label，因此總數(shù)為 mDims.n()

mFileBatch.resize(mDims.n()*mImageSize, 0);

mFileLabels.resize(mDims.n(), 0);

reset(0);

}

// reset操作

void reset(int firstBatch)

{

mBatchCount = 0;

mFileCount = 0;

mFileBatchPos = mDims.n();

skip(firstBatch);

}

/**

* stream.next()每調(diào)用一次，就使用batch file中的數(shù)據(jù)(讀取后首先是變量名為mFileBatch的buffer)填充一個mBatch

* @return 是否填充成功

bool next()

{

//已經(jīng)讀取到最大批數(shù)量了，返回false

if (mBatchCount == mMaxBatches)

return false;

// 將mFileBatch（相當(dāng)于buffer）中的內(nèi)容拷貝到mBatch中，

//由于mFileBatch和mBatch大小有可能不一樣,所以才這么寫

for (int csize = 1, batchPos = 0; batchPos < mBatchSize; batchPos += csize, mFileBatchPos += csize)

{

assert(mFileBatchPos > 0 && mFileBatchPos <= mDims.n());

//調(diào)用update函數(shù)，讀取batches文件夾中的 batch{i} 文件，讀取失敗的話直接在這里返回false，

//調(diào)用update函數(shù)會使 mFileBatchPos=0，這是合理的，因為還沒有開始往 mBatch 拷貝數(shù)據(jù)

if (mFileBatchPos == mDims.n() && !update())

return false;

//一次從batch文件中讀取 csize 張圖片，

//由于mFileBatch和mBatch大小有可能不一樣所以借助 mFileBatchPos 和 batchpos 來指示batch文件和mbatch中的當(dāng)前操作(讀取或存儲)位置

//所以csize取二者之間較小值

// copy the smaller of: elements left to fulfill the request, or elements left in the file buffer.

csize = std::min(mBatchSize - batchPos, mDims.n() - mFileBatchPos);

//將 mFileBatch 和 mFileLabels 中存放的batch文件的內(nèi)容復(fù)制到 mBatch 和 mLabels 中

std::copy_n(getFileBatch() + mFileBatchPos * mImageSize, csize * mImageSize, getBatch() + batchPos * mImageSize);

std::copy_n(getFileLabels() + mFileBatchPos, csize, getLabels() + batchPos);

}

// mBatchCount自增，指示當(dāng)前填充了多少個mBatch

mBatchCount++;

return true;

}

/**

* 跳過前面多少個batch

* @param skipCount 跳過的batch的個數(shù)

void skip(int skipCount)

{

//如果mBatchSize 大于等于 mDims.n()，并且 mBatchSize%mDims.n() == 0，

//換句話說batchsteam中的batchsize(比如100)，比batch{i}文件的batchsize(比如50)大，并且能整除.

//那么batchstream中一個 batch，相當(dāng)于 mBatchSize / mDims.n()個batch 個batch{i}文件

//舉個例子：batchsteam中batchsize=100，batch{i}文件中batchsize=50，那么batchsteam中一個batch相當(dāng)于兩個batch{i}文件

//那么在batchstream中跳過一個 batch，相當(dāng)于跳過 mBatchSize / mDims.n() 個 batch{i}文件

//所以才有 mFileCount += skipCount * mBatchSize / mDims.n();

//這時直接通過修改mFileCount的數(shù)值來讀取剩下的batch文件

if (mBatchSize >= mDims.n() && mBatchSize%mDims.n() == 0 && mFileBatchPos == mDims.n())

{

mFileCount += skipCount * mBatchSize / mDims.n();

return;

}

//其他情況：batchsteam中的batchsize不能整除batch{i}文件的batchsize

//循環(huán)調(diào)用 next() 讀取batch{i}文件，讀取skipCount個，由于next() 會改變 mBatchCount 的值，所以先暫存，再取出

int x = mBatchCount;

for (int i = 0; i < skipCount; i++)

next();

mBatchCount = x;

}

//獲取batchsteam中的 batch 和 label 的首地址， batch文件中的內(nèi)容讀取后首先是放在 mFileBatch 和 mFileLabels 中，

//但最終會被復(fù)制到 mBatch和mLabels中，校準(zhǔn)使用的就是 mBatch 和mLabels，而不是直接從batch file中讀取進(jìn)來的mFileBatch和mFileLabels

float *getBatch() { return &mBatch[0]; }

float *getLabels() { return &mLabels[0]; }

//mBatchCount表示填充了多少個 mBatch 的數(shù)量

//mBatchSize表示填充mBatch時使用的batchsize

int getBatchesRead() const { return mBatchCount; }

int getBatchSize() const { return mBatchSize; }

//獲取圖片的shape信息，這個在mBatch和mFileBatch中是一樣的

nvinfer1::DimsNCHW getDims() const { return mDims; }

private:

//batch文件（如batch0）中的圖像數(shù)據(jù)和標(biāo)簽數(shù)據(jù)存放在 mFileBatch 和 mFileLabels 中，此處返回他們的地址

float* getFileBatch() { return &mFileBatch[0]; }

float* getFileLabels() { return &mFileLabels[0]; }

//此函數(shù)用于依次讀取 batches文件夾下的 batch{i} 文件，并將讀取到的內(nèi)容存放在mFileBatch和mFileLabels中，讀取成功返回true，否則返回false

bool update()

{

//依次讀取 batches文件夾下的 batch{i} 文件，mFileCount變量自增，指向下一個batch文件也就是 batch{i+1} 文件

std::string inputFileName = locateFile(std::string("batches/batch") + std::to_string(mFileCount++));

FILE * file = fopen(inputFileName.c_str(), "rb");

if (!file)

return false;

//從batch文件讀取當(dāng)前 batch 的 shape 信息（圖像數(shù)據(jù)的shape）

int d[4];

fread(d, sizeof(int), 4, file);

assert(mDims.n() == d[0] && mDims.c() == d[1] && mDims.h() == d[2] && mDims.w() == d[3]);

//從batch文件讀取圖像數(shù)據(jù)（精度為float，大小為mDims.n()*mImageSize ），存放到 mFileBatch 中

//從batch文件讀取標(biāo)簽數(shù)據(jù)（精度為float，大小為mDims.n()），存放到mFileLabels中

size_t readInputCount = fread(getFileBatch(), sizeof(float), mDims.n()*mImageSize, file);

size_t readLabelCount = fread(getFileLabels(), sizeof(float), mDims.n(), file);;

assert(readInputCount == size_t(mDims.n()*mImageSize) && readLabelCount == size_t(mDims.n()));

fclose(file);

//每讀取一個batch文件，mFileBatchPos置零，也就是說新讀取的batch文件內(nèi)容 mFileBatch 還沒有開始往 mBatch 拷貝

mFileBatchPos = 0;

//讀取成功返回true

return true;

}

//stream中的批尺寸和最大批數(shù)量，每填充一個mBatch，mBatchCount 自增1

int mBatchSize{ 0 };

int mMaxBatches{ 0 };

int mBatchCount{ 0 };

//mFileCount指向batches文件夾中的batch文件，就跟指針一樣，讀完一個batch，自增1

//mFileBatchPos在一個batch中當(dāng)前操作的位置

int mFileCount{ 0 }, mFileBatchPos{ 0 };

//batchstream中的圖片大小，一般要求跟batch文件中的大小一致，初值為0

int mImageSize{ 0 };

//batch文件中的數(shù)據(jù)的shape

nvinfer1::DimsNCHW mDims;

// 從 batch文件中讀到的圖像數(shù)據(jù)和標(biāo)簽數(shù)據(jù)最終要放到這里來，這個是最終校準(zhǔn)時使用的

std::vector<float> mBatch;

std::vector<float> mLabels;

//用以存取從 batch文件中讀到的圖像數(shù)據(jù)和標(biāo)簽數(shù)據(jù)，相當(dāng)于buffer

std::vector<float> mFileBatch;

std::vector<float> mFileLabels;

};

#endif

7 結(jié)果

myself@admin:~/workspace/study/tensorrt/bin$ ./sample_int8 mnist

FP32 run:400 batches of size 100 starting at 100

........................................

Top1: 0.9904, Top5: 1

Processing 40000 images averaged 0.00167893 ms/image and 0.167893 ms/batch.

FP16 run:400 batches of size 100 starting at 100

Engine could not be created at this precision

INT8 run:400 batches of size 100 starting at 100

........................................

Top1: 0.9908, Top5: 1

Processing 40000 images averaged 0.0013438 ms/image and 0.13438 ms/batch.

從這例程中也忽然發(fā)現(xiàn)在TensorRT中 1080ti GPU竟然不支持 FP16 mode，雖然1080ti官方的參數(shù)上是支持 float16的，但是在TensorRT中竟然不能使用。查了一下，是因為 1080ti的float16 吞吐量太低（throughput），效率太低，應(yīng)該是TensorRT對float16也進(jìn)行了條件限制，吞吐量太低的不支持。

從資料中得知，只有 Tesla P100, Quadro GP100, and Jetson TX1/TX2 支持 full-rate FP16 performance，應(yīng)該也就只有這些才支持 TensorRT的FP16吧。新出的 TITAN V 加了tensor core，float16半精度性能有很大提升，應(yīng)該也支持？不過有意思的是jetson TX1和 TX2 卻能支持 FP16，反而不支持INT8.

可以參考下面資料：

FP16 –half=true option doesn’t work on GTX 1080 TI although it runs ./sample_int8 INT8
FP16 support on gtx 1060 and 1080

The only GPUs with full-rate FP16 performance are Tesla P100, Quadro GP100, and Jetson TX1/TX2. All GPUs with compute capability 6.1 (e.g. GTX 1050, 1060, 1070, 1080, Pascal Titan X, Titan Xp, Tesla P40, etc.) have low-rate FP16 performance. It’s not the fast path on these GPUs. All of these GPUs should support “full rate” INT8 performance, however.

從結(jié)果上看：

INT8 MODE：Top 1 0.9908，速度：0.0013438 ms/image ；

FP32 MODE : Top 1 0.9904，速度：0.00167893 ms/image；

準(zhǔn)確率竟然還高那么一點點，速度上大概快了20%。

參考

TensorRT Developer Guide

cutoff and quantile parameters in TensorRT

FP16 –half=true option doesn’t work on GTX 1080 TI although it runs ./sample_int8 INT8

FP16 support on gtx 1060 and 1080

總結(jié)

以上是生活随笔為你收集整理的TensorRT(6)-INT8 inference的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯，歡迎將生活随笔推薦給好友。

上一篇： TensorRT(5)-INT8校准原理
下一篇： YoloV5代码详细解读