-
Notifications
You must be signed in to change notification settings - Fork 3
/
Copy pathUser Guide.html
281 lines (173 loc) · 15 KB
/
User Guide.html
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
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
<!DOCTYPE html>
<html>
<head>
<meta charset="utf-8">
<meta name="viewport" content="width=device-width, initial-scale=1.0">
<title>GIE User Guide</title>
<link rel="stylesheet" href="https://stackedit.io/res-min/themes/base.css" />
<script type="text/javascript" src="https://cdn.mathjax.org/mathjax/latest/MathJax.js?config=TeX-AMS_HTML"></script>
</head>
<body><div class="container"><h1 id="gie-user-guide">GIE User Guide</h1>
<p>The NVIDIA GPU Inference Engine (GIE) is a C++ library that facilitates high performance inference on NVIDIA GPUs. It takes a network definition and optimizes it by merging tensors and layers, transforming weights, choosing efficient intermediate data formats, and selecting from a large kernel catalog based on layer parameters and measured performance.</p>
<p>This release of GIE is built with gcc 4.8.</p>
<p>GIE has the following layer types:</p>
<ul>
<li><strong>Convolution</strong>, with or without bias. Currently only 2D convolutions (i.e. 4D input and output tensors) are supported. <strong>Note</strong>: The operation this layer performs is actually a correlation, which is a consideration if you are formatting weights to import via GIE’s API rather than the caffe parser library. </li>
<li><strong>Activation</strong>: ReLU, tanh and sigmoid.</li>
<li><strong>Pooling</strong>: max and average.</li>
<li><strong>Scale</strong>: per-tensor, per channel or per-weight affine transformation and exponentiation by constant values. <strong>Batch Normalization</strong> can be implemented using the Scale layer.</li>
<li><strong>ElementWise</strong>: sum, product or max of two tensors.</li>
<li><strong>LRN</strong>: cross-channel only.</li>
<li><strong>Fully-connected</strong> with or without bias</li>
<li><strong>SoftMax</strong>: cross-channel only</li>
<li><strong>Deconvolution</strong>, with and without bias</li>
</ul>
<p>While GIE is independent of any framework, the package includes a parser for caffe models named NvCaffeParser, which provides a simple mechanism for importing network definitions. NvCaffeParser uses the above layers to implement caffe’s Convolution, ReLU, Sigmoid, TanH, Pooling, Power, BatchNorm, Eltwise, LRN, InnerProduct, SoftMax, Scale, and Deconvolution layers. Caffe features not supported by GIE include:</p>
<ul>
<li>Deconvolution groups</li>
<li>PReLU</li>
<li>Scale, other than per-channel scaling</li>
<li><p>EltWise with more than two inputs</p>
<p><strong>Note:</strong> GIE’s caffe parser does not support legacy formats in caffe prototxt - in particular, layer types are expected to be expressed in the prototxt as strings delimited by double quotes.</p></li>
</ul>
<h1 id="getting-started">Getting Started</h1>
<p>There are two phases to using GIE:</p>
<ul>
<li>In the <em>build</em> phase, the toolkit takes a network definition, performs optimizations, and generates the inference engine. </li>
<li>In the <em>execution</em> phase, the engine runs inference tasks using input and output buffers on the GPU.</li>
</ul>
<p>The build phase can take considerable time, especially when running on embedded platforms. So a typical application will build an engine once, and then serialize it for later use.</p>
<p>The build phase performs the following optimizations on the layer graph:</p>
<ul>
<li>elimination of layers whose outputs are not used</li>
<li>fusion of convolution, bias and ReLU operations</li>
<li>aggregation of operations with sufficiently similar parameters and the same source tensor (for example, the 1x1 convolutions in GoogleNet v5’s inception module)</li>
<li>elision of concatenation layers by directing layer outputs to the correct eventual destination.</li>
</ul>
<p>In addition it runs layers on dummy data to select the fastest from its kernel catalog, and performs weight preformatting and memory optimization where appropriate.</p>
<h2 id="the-network-definition">The Network Definition</h2>
<p>A <em>network definition</em> consists of a sequence of layers, and a set of tensors. </p>
<p>Each <em>layer</em> computes a set of output tensors from a set of input tensors. Layers have parameters, e.g. convolution size and stride, and convolution filter weights.</p>
<p>A <em>tensor</em> is either an input to the network, or an output of a layer. Tensors have a datatype specifying their precision (16- and 32-bit floats are currently supported) and three dimensions (channels, width, height). The dimensions of an input tensor are defined by the application, and for output tensors they are inferred by the builder.</p>
<p>Each layer and tensor has a name, which can be useful when profiling or reading GIE’s build log (see Logging below) </p>
<p>When using the caffe parser, tensor and layer names are taken from the caffe prototxt. </p>
<h2 id="samplemnist">SampleMNIST</h2>
<p>The MNIST sample demonstrates typical build and execution phases using a caffe model trained on the MNIST dataset using the <a href="https://github.com/NVIDIA/DIGITS/blob/master/docs/GettingStarted.md">NVIDIA DIGITS tutorial</a>.</p>
<h3 id="logging">Logging</h3>
<p>GIE requires a logging interface to be implemented, through which it reports errors, warnings, and informational messages. </p>
<pre><code>class Logger : public ILogger
{
void log(Severity severity, const char* msg) override
{
// suppress info-level messages
if (severity != Severity::kINFO)
std::cout << msg << std::endl;
}
} gLogger;
</code></pre>
<p>Here we suppress informational messages, and report only warnings and errors.</p>
<h3 id="the-build-phase-caffetogiemodel">The Build Phase - <code>caffeToGIEModel</code></h3>
<p>First we create the GIE builder. The application must implement a logging interface, through which GIE will provide information about optimization stages during the build phase, and also warnings and error information. </p>
<pre><code>IBuilder* builder = createInferBuilder(gLogger);
</code></pre>
<p>Then we create the network definition structure, which we populate from a caffe model using the caffe parser library: </p>
<pre><code>INetworkDefinition* network = infer->createNetwork();
CaffeParser* parser = createCaffeParser();
std::unordered_map<std::string, infer1::Tensor> blobNameToTensor;
const IBlobNameToTensor* blobNameToTensor =
parser->parse(locateFile(deployFile).c_str(),
locateFile(modelFile).c_str(),
*network,
DataType::kFLOAT);
</code></pre>
<p>In this sample, we instruct the parser to generate a network whose weights are 32-bit floats. As well as populating the network definition, the parser returns a dictionary that maps from caffe blob names to GIE tensors. </p>
<p><strong>Note:</strong> A GIE network definition has no notion of in-place operation - e.g. the input and output tensors of a ReLU are different. When a caffe network uses an in-place operation, the GIE tensor returned in the dictionary corresponds to the last write to that blob. For example, if a convolution creates a blob and is followed by an in-place ReLU, that blob’s name will map to the GIE tensor which is the output of the ReLU.</p>
<p>Since the caffe model does not tell us which tensors are the outputs of the network, we need to specify these explicitly after parsing:</p>
<pre><code>for (auto& s : outputs)
network->markOutput(*blobNameToTensor->find(s.c_str()));
</code></pre>
<p>There is no restriction on the number of output tensors, but marking a tensor as an output may prohibit some optimizations on that tensor.</p>
<p><strong>Note:</strong> at this point we have parsed the caffe model to obtain the network definition, and are ready to create the engine. However, we cannot yet release the parser object because the network definition holds weights by reference into the caffe model, not by value. It is only during the build process that the weights are read from the caffemodel.</p>
<p>We next build the engine from the network definition:</p>
<pre><code>builder->setMaxBatchSize(maxBatchSize);
builder->setMaxWorkspaceSize(1 << 20);
ICudaEngine* engine = builder->buildCudaEngine(*network);
</code></pre>
<ul>
<li><code>maxBatchSize</code> is the size for which the engine will be tuned. At execution time, smaller batches may be used, but not larger. Note that execution of smaller batch sizes may be slower than with a GIE engine optimized for that size.</li>
<li><code>maxWorkspaceSize</code> is the maximum amount of scratch space which the engine may use at runtime</li>
</ul>
<p>We could use the engine directly, but here we serialize it to a stream, which is the typical usage mode for GIE:</p>
<pre><code>engine->serialize(gieModelStream);
</code></pre>
<h3 id="deserializing-the-engine">Deserializing the engine</h3>
<p>To deserialize the engine, we create a GIE runtime object:</p>
<pre><code>IRuntime* runtime = createInferRuntime(gLogger);
ICudaEngine* engine = runtime->deserializeCudaEngine(gieModelStream);
</code></pre>
<p>We also need to create an execution context. One engine can support multiple contexts, allowing inference to be performed on multiple batches simultaneously while sharing the same weights.</p>
<pre><code>IExecutionContext *context = engine->createExecutionContext();
</code></pre>
<p><strong>Note:</strong> Serialized engines are not portable across platforms or GIE versions.</p>
<h3 id="the-execution-phase-doinference">The Execution Phase - <code>doInference()</code></h3>
<p>The input to the engine is an array of pointers to input and output buffers on the GPU (<strong>Note:</strong> All GIE inputs and outputs are in contiguous NCHW format.) The engine can be queried for the buffer indices, using the tensor names assigned when the network was created. </p>
<pre><code>int inputIndex = engine->getBindingIndex(INPUT_BLOB_NAME),
outputIndex = engine->getBindingIndex(OUTPUT_BLOB_NAME);
</code></pre>
<p>In a typical production case, GIE will execute asynchronously. The <code>enqueue()</code> method will add add kernels to a cuda stream specified by the application, which may then wait on that stream for completion. The fourth parameter to <code>enqueue()</code> is an optional <code>cudaEvent</code> which will be signaled when the input buffers are no longer in use and can be refilled. </p>
<p>In this sample we simply copy the input buffer to the GPU, run inference, then copy the result back and wait on the stream:</p>
<pre><code>cudaMemcpyAsync(<...>, cudaMemcpyHostToDevice, stream);
context.enqueue(batchSize, buffers, stream, nullptr);
cudaMemcpyAsync(<...>, cudaMemcpyDeviceToHost, stream);
cudaStreamSynchronize(stream);
</code></pre>
<p><strong>Note:</strong> The batch size must be at most the value specified when the engine was created.</p>
<h2 id="samplegooglenet">SampleGoogleNet</h2>
<p>SampleGoogleNet illustrates layer-based profiling and GIE’s half2 mode, which is the fastest mode for batch sizes greater than one on platforms which natively support 16-bit inference . </p>
<h3 id="profiling">Profiling</h3>
<p>To profile a network, implement the <code>IProfiler</code> interface and add the profiler to the execution context:</p>
<pre><code>context.profiler = &gProfiler;
</code></pre>
<p>Profiling is not currently supported for asynchronous execution, and so we must use GIE’s synchronous <code>execute()</code> method:</p>
<pre><code>for (int i = 0; i < TIMING_ITERATIONS;i++)
engine->execute(context, buffers);
</code></pre>
<p>After execution has completed, the profiler callback is called once for every layer. The sample accumulates layer times over invocations, and averages the time for each layer at the end. </p>
<p>Observe that the layer names are modified by GIE’s layer-combining operations.</p>
<h3 id="half2-mode">half2 mode</h3>
<p>GIE can use 16-bit instead of 32-bit arithmetic and tensors, but this alone may not deliver significant performance benefits. Half2 is an execution mode where internal tensors interleave 16 bits from adjacent pairs of images, and is the fastest mode of operation for batch sizes greater than one.</p>
<p>To use half2 mode, two additional steps are required:</p>
<p>1) create an input network with 16-bit weights, by supplying the <code>DataType::kHALF2</code> parameter to the parser </p>
<pre><code>const IBlobNameToTensor *blobNameToTensor =
parser->parse(locateFile(deployFile).c_str(),
locateFile(modelFile).c_str(),
*network,
DataType::kHALF);
</code></pre>
<p>2) set the ‘half2mode’ mode in the builder when building the engine:</p>
<pre><code>builder->setHalf2Mode(true);
</code></pre>
<h2 id="using-gie-with-multiple-gpus">Using GIE with multiple GPUs</h2>
<p>Each <code>ICudaEngine</code> object is bound to a specific GPU when it is instantiated, either by the builder or on deserialization. To select the GPU, use <code>cudaSetDevice()</code> before calling the builder or deserializing the engine. Each <code>IExecutionContext</code> is bound to the same GPU as the engine from which it was created. When calling <code>execute()</code> or <code>enqueue()</code>, ensure that the thread is associated with the correct device by calling <code>cudaSetDevice()</code> if necessary.</p>
<h2 id="data-formats">Data Formats</h2>
<p>GIE network inputs and outputs are 32-bit tensors in contiguous NCHW format.</p>
<p>For weights:</p>
<ul>
<li>Convolution weights are in contiguous KCRS format, where K indexes over output channels, C over input channels, and R and S over the height and width of the convolution, respectively.</li>
<li>Fully Connected weights are in contiguous row-major layout</li>
<li>Deconvolution weights are in contiguous CKRS format (where C, K, R and S are as above).</li>
</ul>
<h1 id="faq">FAQ</h1>
<p><strong>Q: How can I use my own layer types in conjunction with GIE?</strong> <br>
A: This release of GIE doesn’t support custom layers. To use your own layer in the middle of GIE, create two GIE pipelines, one to run before your layer and one to run afterwards.</p>
<pre><code>IExecutionContext *contextA = engineA->createExecutionContext();
IExecutionContext *contextB = engineB->createExecutionContext();
<...>
contextA.enqueue(batchSize, buffersA, stream, nullptr);
myLayer(outputFromA, inputToB, stream);
contextB.enqueue(batchSize, buffersB, stream, nullptr);
</code></pre>
<p>The first GIE pipeline will read the input and write to <code>outputFromA</code>, and the second will read from <code>inputToB</code> and generate the final output.</p>
<p><strong>Q: How do I create an engine optimized for several possible batch sizes?</strong> <br>
A: While GIE allows an engine optimized for a given batch size to run at any smaller size, the performance for those smaller sizes may not be as well-optimized. To optimize for multiple different batch sizes, run the builder and serialize an engine for each batch size. A future release of GIE will be able to optimize a single engine for multiple batch sizes, thereby allowing for sharing of weights where layers at different batch sizes use the same weight formats. </p></div></body>
</html>