Padding Mask

Figure 13. cuDNN graph depicting DAG:Padding Mask

Causal Mask

Figure 14. cuDNN graph depicting DAG:Causal Mask

Softmax

Figure 15. cuDNN graph depicting DAG:Softmax

Dropout

Figure 16. cuDNN graph depicting DAG:Dropout

g₄ is capable of storing an intermediate tensor to global memory marked as S, which can be used for pattern 7. Both DAG:Softmax and DAG:Dropout have this capability. You should set S as the output from the last DAG in the graph.

The tensor descriptor marked as S needs to have the CUDNN_ATTR_TENSOR_REORDERING_MODE set to CUDNN_TENSOR_REORDERING_F16x16. This is because the tensor is stored in a special format and can only be consumed by pattern 7.

g₅, g₆, and g₇ can only support a fixed DAG. We are working towards generalizing these graphs.

Figure 17. cuDNN graph depicting g₅

g6 represents the backward pass of Softmax and masking, to get dP.

Figure 18. cuDNN graph depicting g₆

g7 is the transpose of dP the output of g6.

Figure 19. cuDNN graph depicting g₇

For Convolutions

// Set the math type to allow cuDNN to use Tensor Cores:
checkCudnnErr(cudnnSetConvolutionMathType(cudnnConvDesc, CUDNN_TENSOR_OP_MATH_ALLOW_CONVERSION));

For RNNs

// Set the math type to allow cuDNN to use Tensor Cores:
checkCudnnErr(cudnnSetRNNMatrixMathType(cudnnRnnDesc, CUDNN_TENSOR_OP_MATH_ALLOW_CONVERSION));

Convolution (convolution mode set to CUDNN_CROSS_CORRELATION)

$y_{\mathit{n,\; k,\; p,\; q}}=\sum _c^C\sum _r^R\sum _s^S{x}_{\mathit{n,\; c,\; p+r,\; q+s}}\times w_{\mathit{k,c,r,s}}$

Convolution with Padding

$x_{\mathit{<0,\; <0}}=0$

$x_{\mathit{>H,\; >W}}=0$

$y_{\mathit{n,\; k,\; p,\; q}}=\sum _c^C\sum _r^R\sum _s^S{x}_{\mathit{n,\; c,\; p+r-pad,\; q+s-pad}}\times w_{\mathit{k,c,r,s}}$

Convolution with Subsample-Striding

$y_{\mathit{n,\; k,\; p,\; q}}=\sum _c^C\sum _r^R\sum _s^S{x}_{\mathit{n,\; c,\; (p*u)\; +\; r,\; (q*v)\; +\; s}}\times w_{\mathit{k,c,r,s}}$

Convolution with Dilation

$y_{\mathit{n,\; k,\; p,\; q}}=\sum _c^C\sum _r^R\sum _s^S{x}_{\mathit{n,\; c,\; p\; +\; (r*dilh),\; q\; +\; (s*dilw)}}\times w_{\mathit{k,c,r,s}}$

Convolution (convolution mode set to CUDNN_CONVOLUTION)

$y_{\mathit{n,\; k,\; p,\; q}}=\sum _c^C\sum _r^R\sum _s^S{x}_{\mathit{n,\; c,\; p\; +\; r,\; q\; +\; s}}\times w_{\mathit{k,\; c,\; R-r-1,\; S-s-1}}$

Convolution using Grouped Convolution

$C_g=\frac{C}{G}$

$K_g=\frac{K}{G}$

$y_{\mathit{n,\; k,\; p,\; q}}=\sum _c^{C_g}\sum _r^R\sum _s^S{x}_{\mathit{n,\; Cg*floor(k/Kg)+c,\; p+r,\; q+s}}\times w_{\mathit{k,c,r,s}}$

Basic Idea

Conceptually, in grouped convolutions, the input channels and the filter channels are split into a groupCount number of independent groups, with each group having a reduced number of channels. The convolution operation is then performed separately on these input and filter groups.

For example, consider the following: if the number of input channels is 4, and the number of filter channels of 12. For a normal, ungrouped convolution, the number of computation operations performed are 12*4.

If the groupCount is set to 2, then there are now two input channel groups of two input channels each, and two filter channel groups of six filter channels each.

As a result, each grouped convolution will now perform 2*6 computation operations, and two such grouped convolutions are performed. Hence the computation savings are 2x: (12*4)/(2*(2*6)) .

cuDNN Grouped Convolution

Example