Is there somewhere I can find more detail about this? You’re right about launching a few kernels. Each batch launches a kernel for the im2col computation and a kernel for the gemm (done through cuBLAS). However, when I profile the loop, I see that most of the launch time is taken up by the initial im2col call (i.e. in the first batch). Synchronizing with cudaDeviceSynchronize() after each call produces these times:
batch: 0
im2col: 0.000030s
gemm: 0.000026s
batch: 1
im2col: 0.000012s
gemm: 0.000008s
batch: 2
im2col: 0.000008s
gemm: 0.000007s
batch: 3
im2col: 0.000008s
gemm: 0.000007s
batch: 4
im2col: 0.000008s
gemm: 0.000007s
batch: 5
im2col: 0.000008s
gemm: 0.000007s

I’d presumed the speed up for the later batches was due to caching the columns (for im2col) and the weights (for gemm), but I don’t know how I’d do this with a single kernel. Or this this something that cudnn makes possible? As in, does cudnnConvolutionForward 7 implement the entire operation on a single kernel?