Fermi-class and later GPUs are capable of concurrently running
multiple kernels, provided they are launched in different streams and
with block sizes that are small enough that a single kernel will not
fill the whole GPU. The code of Listing 6-5 (Lines 9-14) will cause
kernels to run concurrently, provided the number of blocks in each
kernel launch is small enough. Since the kernels can only communicate
through global memory, we can add some instrumentation to AddKernel() to
track how many kernels are running concurrently. Using a “kernel
concurrency tracking” structure:
static const int g_maxStreams = 8;
typedef struct KernelConcurrencyData_st {
int mask; // mask of active kernels
int maskMax; // atomic max of mask popcount
int masks[g_maxStreams];
int count; // number of active kernels
int countMax; // atomic max of kernel count
int counts[g_maxStreams];
} KernelConcurrencyData;
we can add code to AddKernel() to “check in” and “check out” at the
beginning and end of the function. The “check in” takes the “kernel id”
parameter kid (a value in the range 0..NumStreams-1 passed
to the kernel), OR’s a mask corresponding to the kernel ID into a
global, and records that kernel’s view of the global. Note that
atomicOr() returns the value that was in the memory location before the
OR was performed.
Similarly, this code tracks the number of active kernels by
incrementing kernelData-\>count and calling atomicMax()
on a shared global.
// check in, and record active kernel mask and count
// as seen by this kernel.
if ( kernelData && blockIdx.x==0 && threadIdx.x == 0 ) {
int myMask = atomicOr( &kernelData->mask, 1<<kid );
kernelData->masks[kid] = myMask | (1<<kid);
int myCount = atomicAdd( &kernelData->count, 1 );
atomicMax( &kernelData->countMax, myCount+1 );
kernelData->counts[kid] = myCount+1;
}
At the bottom of the kernel, similar code clears the mask and decrements the active-kernel count:
// check out
if ( kernelData && blockIdx.x==0 && threadIdx.x==0 ) {
atomicAnd( &kernelData->mask, ~(1<<kid) );
atomicAdd( &kernelData->count, -1 );
}
The kernelData parameter refers to a __device__ variable declared at
file scope:
__device__ KernelConcurrencyData g_kernelData;
Remember, the pointer to g_kernelData must be obtained
by calling cudaGetSymbolAddress(). It is possible to write code that
references &g_kernelData, but CUDA’s language
integration will not correctly resolve the address.
The concurrencyKernelKernel.cu program adds support for a command
line option blocksPerSM to specify the number of blocks
with which to launch these kernels. It will generate a report on the
number of kernels that were active. Two sample invocations of
concurrencyKernelKernel are as follows:
$ ./concurrencyKernelKernel --blocksPerSM 2
Using 2 blocks per SM on GPU with 14 SMs = 28 blocks
Timing sequential operations... Kernel data:
Masks: ( 0x1 0x0 0x0 0x0 0x0 0x0 0x0 0x0 )
Up to 1 kernels were active: (0x1 0x0 0x0 0x0 0x0 0x0 0x0 0x0 )
Timing concurrent operations...
Kernel data:
Masks: ( 0x1 0x3 0x7 0xe 0x1c 0x38 0x60 0xe0 )
Up to 3 kernels were active: (0x1 0x2 0x3 0x3 0x3 0x3 0x2 0x3 )
$ ./concurrencyKernelKernel --blocksPerSM 3
Using 3 blocks per SM on GPU with 14 SMs = 42 blocks
Timing sequential operations... Kernel data:
Masks: ( 0x1 0x0 0x0 0x0 0x0 0x0 0x0 0x0 )
Up to 1 kernels were active: (0x1 0x0 0x0 0x0 0x0 0x0 0x0 0x0 )
Timing concurrent operations... Kernel data:
Masks: ( 0x1 0x3 0x6 0xc 0x10 0x30 0x60 0x80 )
Up to 2 kernels were active: (0x1 0x2 0x2 0x2 0x1 0x2 0x2 0x1 )Note that blocksPerSM is the number of blocks specified
to each kernel launch, so a total of
numStreams\*blocksPerSM blocks are launched in
numStreams separate kernels. You can see that the hardware
can run more kernels concurrently when the kernel grids are smaller, but
there is no performance benefit to concurrent kernel processing for the
workload discussed in this chapter.