Ядро lauch, указывающее поток, но с размером общей памяти по умолчанию

Мне нужно указать поток для запуска ядра в CUDA. Ядро использует некоторую разделяемую память, размер которой определен в коде ядра.

static const int cBlockSize = 256;

__global__ fooKernel(void* param)
{
__shared__ uint32_t words[cBlockSize/16];
// implementation follows, using 2 bits of shared memory per thread
}

Однако параметр размера разделяемой памяти идет перед параметром потока в выражении запуска ядра. Так как же сказать CUDA использовать размер разделяемой памяти, заданный кодом ядра, и игнорировать то, что в коде запуска?

fooKernel<<<N/cBlockSize, cBlockSize, /* What to put here? */, stream>>>(param);

Очевидно, я бы хотел избежать дублирования кода (cBlockSize/16)*sizeof(uint32_t) там снова. На самом деле выражение является более сложным.