Dynamically allocated shared memory in CUDA. Execution Configuration

Question

What does Nvidia mean by this?

Ns is of type size_t and specifies the number of bytes in shared memory that is dynamically allocated per block for this call in addition to the statically allocated memory; this dynamically allocated memory is used by any of the variables declared as an external array as mentioned in __shared__; Ns is an optional argument which defaults to 0;

Size of shared memory in my GPU is 48kB. For example I want to run a kernel with 4 parallel blocks, every one of them uses 12kB of shared memory.

In order to do that, should I launch the kernel this way

kernel<<< gridSize, blockSize, 12 * 1024 >>>();

or should the third argument be 48 * 1024?

The third argument specifies how much dynamically allocated shared memory is used by the kernel, so the third argument should be 12 * 1024. But if you want to run four kernels in parallel, I expect you will be using streams, so there should be also a fourth argument in your kernel calls. — Vitality
– Vitality, Commented Feb 25, 2015 at 20:19
Thanks! By the way I meant not '4 kernels at the same time', but '1 kernel on 4 threadblocks at the same time'. I guess I still need to use kernel<<< gridSize, blockSize, 12 * 1024 * 1024 >>>(); — artur me
– artur me, Commented Feb 26, 2015 at 1:15

Christian Sarofeen · Accepted Answer · 2015-02-25 21:39:31Z

1

Ns in a size in bytes. If you want to reserve 12kB of shared memory you would do 12*1024*1024.

I doubt you want to do this. Ns value is PER BLOCK. So it is the amount of shared memory per block executing on the device. I'm guessing you'd like to do something around the lines of 12*1024*1024/number_of_blocks;

Kernel launching with concurrency: If as mentioned in a comment, you are using streams there is a fourth input in the kernel launch which is the cuda stream.

If you want to launch a kernel on another stream without any shared memory it will look like:

kernel_name<<<128, 128, 0, mystream>>>(...);

but concurrency is a whole different issue.

answered Feb 25, 2015 at 21:39

Christian Sarofeen

2,28013 silver badges18 bronze badges

Sign up to request clarification or add additional context in comments.

1 Comment

artur me Over a year ago

Yep, you're a right, I wanted 12 * 1024 * 1024, It was kind of misprint. Thank you for your answer.

Collectives™ on Stack Overflow

Dynamically allocated shared memory in CUDA. Execution Configuration

1 Answer 1

1 Comment

Your Answer

Hot Network Questions

Collectives™ on Stack Overflow

1 Answer 1

1 Comment

Your Answer

Sign up or log in

Post as a guest

Related