What are templated methods in C++ / CUDA with 3 angle brackets (<<<)?

Question

Forgive what may be a simple question, my C++ is rusty. I'm working on a machine learning application in C that uses CUDA for some calculations, and I've found the following line of code.

Just curious how to parse this. It looks like a templated method, but I dont understand the triple angle-brackets '<<<'. What's going on here?

backward_scale_kernel<<<n, BLOCK>>>(x_norm, delta, batch, n, size, scale_updates);

For context "n" is passed in as a function parameter, and I can't find where BLOCK is defined or assigned.

Answer 1

I'm sure there is a good duplicate for this already, but the <<< >>> decorator you are seeing is a CUDA runtime API syntax extension which allows the execution parameters of a CUDA kernel call to be specified.

The full syntax is

kernel_function<<<grid dimensions, block dimensions, dynamic shared memory, stream ID>>>( ....arguments....)

The CUDA runtime API front end expands this syntax into a pair of inline function calls to compiler emitted boilerplate to allow the underlying GPU kernel to be launched at runtime.