CUDA: Difference between revisions

Line 89:

* that the total number of threads not exceed some limit ''t'' (likely bounding the divergence-tracking stacks), and

* that the total number of blocks not exceed some limit ''b'' (likely bounding the warp-scheduling complexity).

A given SM, then, supports '''T''' values through the minimum of {''r''/'''Thrreg''', ''s''/'''~~Thr~~shmem''', and ''t''}; as ~~each thread~~ requires fewer registers and less shared memory, the upper bound converges to ''t''. Motivations for larger blocks include:

A given SM, then, supports '''T''' values through the minimum of {''r''/'''Thrreg''', ''s''/'''Blkshmem''', and ''t''}; as the block requires fewer registers and less shared memory, the upper bound converges to ''t''. Motivations for larger blocks include:

* freedom in the ''b'' dimension exposes parallelism until ''t'' <= ''b'' * '''T'''

* larger maximum possible kernels (an absolute limit exists on grid dimensions)

Line 95:

Motivations for smaller blocks include:

* freedom in the ''t'' dimension exposes parallelism until ''t'' >= ''b'' * '''T'''

* freedom in the ''r'' and ''s'' dimensions exposes parallelism until ''r'' >= ''b'' * '''T''' * '''Thrreg''' or ''s'' >= ''b'' * '''~~T''' * '''Thr~~shmem'''.

* freedom in the ''r'' and ''s'' dimensions exposes parallelism until ''r'' >= ''b'' * '''T''' * '''Thrreg''' or ''s'' >= ''b'' * '''Blkshmem'''.

* cheaper per-block operations(?) (<tt>__syncthreads()</tt>, voting, etc)

* support for older hardware and SDKs

@@ Line 89: / Line 89: @@
 * that the total number of threads not exceed some limit ''t'' (likely bounding the divergence-tracking stacks), and
 * that the total number of blocks not exceed some limit ''b'' (likely bounding the warp-scheduling complexity).
-A given SM, then, supports '''T''' values through the minimum of {''r''/'''Thr<sub>reg</sub>''', ''s''/'''Thr<sub>shmem</sub>''', and ''t''}; as each thread requires fewer registers and less shared memory, the upper bound converges to ''t''. Motivations for larger blocks include:
+A given SM, then, supports '''T''' values through the minimum of {''r''/'''Thr<sub>reg</sub>''', ''s''/'''Blk<sub>shmem</sub>''', and ''t''}; as the block requires fewer registers and less shared memory, the upper bound converges to ''t''. Motivations for larger blocks include:
 * freedom in the ''b'' dimension exposes parallelism until ''t'' <= ''b'' * '''T'''
 * larger maximum possible kernels (an absolute limit exists on grid dimensions)
@@ Line 95: / Line 95: @@
 Motivations for smaller blocks include:
 * freedom in the ''t'' dimension exposes parallelism until ''t'' >= ''b'' * '''T'''
-* freedom in the ''r'' and ''s'' dimensions exposes parallelism until ''r'' >= ''b'' * '''T''' * '''Thr<sub>reg</sub>''' or ''s'' >= ''b'' * '''T''' * '''Thr<sub>shmem</sub>'''.
+* freedom in the ''r'' and ''s'' dimensions exposes parallelism until ''r'' >= ''b'' * '''T''' * '''Thr<sub>reg</sub>''' or ''s'' >= ''b'' * '''Blk<sub>shmem</sub>'''.
 * cheaper per-block operations(?) (<tt>__syncthreads()</tt>, voting, etc)
 * support for older hardware and SDKs

Dank (talk | contribs)