dankwiki

CUDA: Difference between revisions

← Older edit

Dank (talk | contribs)

2 user groups
11,325 edits
→Compute Capabilities: resource table
"65K" ?!
 
(18 intermediate revisions by the same user not shown)
Line 1: Line 1:
[[File:Gt200die-big.jpg|right|thumb|A "Fermi" GT200 die]]
[[File:Gt200die-big.jpg|right|thumb|A "Fermi" GT200 die]]
==Hardware==
==Hardware==
NVIDIA maintains a list of [http://www.nvidia.com/object/cuda_learn_products.html supported hardware]. For actual hardware, you'll need the "nvidia.ko" kernel module. Download the <tt>nvidia-kernel-source</tt> and <tt>nvidia-kernel-common</tt> packages, unpack <tt>/usr/src/nvidia-kernel.tar.bz2</tt>, and run <tt>make-kpkg modules_image</tt>. Install the resulting .deb, and modprobe nvidia. You'll see something like this in dmesg output:<pre>nvidia: module license 'NVIDIA' taints kernel.
NVIDIA maintains a list of [http://www.nvidia.com/object/cuda_learn_products.html supported hardware]. You'll need the "nvidia.ko" kernel module. On [[Debian]], use the <tt>nvidia-kernel-dkms</tt> package to build a module appropriate for your kernel (and automatically rebuild it upon kernel upgrades). You can also download the <tt>nvidia-kernel-source</tt> and <tt>nvidia-kernel-common</tt> packages, unpack <tt>/usr/src/nvidia-kernel.tar.bz2</tt>, and run <tt>make-kpkg modules_image</tt>. Install the resulting .deb, and modprobe nvidia. You'll see something like this in dmesg output:<pre>nvidia: module license 'NVIDIA' taints kernel.
Disabling lock debugging due to kernel taint
Disabling lock debugging due to kernel taint
nvidia 0000:07:00.0: enabling device (0000 -> 0003)
nvidia 0000:07:00.0: enabling device (0000 -> 0003)
Line 7: Line 7:
nvidia 0000:07:00.0: setting latency timer to 64
nvidia 0000:07:00.0: setting latency timer to 64
NVRM: loading NVIDIA UNIX x86_64 Kernel Module  190.53  Wed Dec  9 15:29:46 PST 2009</pre>
NVRM: loading NVIDIA UNIX x86_64 Kernel Module  190.53  Wed Dec  9 15:29:46 PST 2009</pre>
Once the module is loaded, CUDA should be able to find the device. See [[CUDA#deviceQuery_Output|below]] for sample outputs. Each device has a '''compute capability''', though this does not encompass all differentiated capabilities (see also <tt>deviceOverlap</tt> and <tt>canMapHostMemory</tt>...). Note that "emulation mode" has been removed as of CUDA Toolkit Version 3.1.
Once the module is loaded, CUDA should be able to find the device. See [[CUDA#deviceQuery_Output|below]] for sample outputs. Each device has a [[CUDA#Compute_Capabilities|compute capability]], though this does not encompass all differentiated capabilities (see also <tt>deviceOverlap</tt> and <tt>canMapHostMemory</tt>...). Note that "emulation mode" has been removed as of CUDA Toolkit Version 3.1.
 
==CUDA model==
==CUDA model==
===Host===
===Host===
Line 31: Line 32:
===Streaming Multiprocessor===
===Streaming Multiprocessor===
* Each SM has a register file, fast local (''shared'') memory, a cache for constant memory, an instruction cache (ROP), a multithreaded instruction dispatcher, and some number of [[#Stream Processor|Stream Processors]] (SPs).
* Each SM has a register file, fast local (''shared'') memory, a cache for constant memory, an instruction cache (ROP), a multithreaded instruction dispatcher, and some number of [[#Stream Processor|Stream Processors]] (SPs).
** 8192 registers for compute capability <= 1.1, otherwise
** 8K registers for compute capability <= 1.1, otherwise
** 16384 for compute capability <= 1.3
** 16K for compute capability <= 1.3, otherwise
** 32K for compute capability <= 2.1, otherwise
** 64K through at least compute capability 3.5
* A group of threads which share a memory and can "synchronize their execution to coördinate accesses to memory" (use a [[barrier]]) form a '''block'''. Each thread has a ''threadId'' within its (three-dimensional) block.
* A group of threads which share a memory and can "synchronize their execution to coördinate accesses to memory" (use a [[barrier]]) form a '''block'''. Each thread has a ''threadId'' within its (three-dimensional) block.
** For a block of dimensions &lt;D<sub>x</sub>, D<sub>y</sub>, D<sub>z</sub>&gt;, the threadId of the thread having index &lt;x, y, z&gt; is (x + y * D<sub>x</sub> + z * D<sub>y</sub> * D<sub>x</sub>).
** For a block of dimensions &lt;D<sub>x</sub>, D<sub>y</sub>, D<sub>z</sub>&gt;, the threadId of the thread having index &lt;x, y, z&gt; is (x + y * D<sub>x</sub> + z * D<sub>y</sub> * D<sub>x</sub>).
* Register allocation is performed per-block, and rounded up to the nearest
* Register allocation is performed per-block, and rounded up to the nearest
** 256 registers per block for compute capability <= 1.1, otherwise
** 256 registers per block for compute capability <= 1.1, otherwise
** 512 registers per block for compute capability <= 1.3.
** 512 registers per block for compute capability <= 1.3
* A group of blocks which share a kernel form a '''grid'''. Each block (and each thread within that block) has a ''blockId'' within its (two-dimensional) grid.
* A group of blocks which share a kernel form a '''grid'''. Each block (and each thread within that block) has a ''blockId'' within its (two-dimensional) grid.
** For a grid of dimensions &lt;D<sub>x</sub>, D<sub>y</sub>&gt;, the blockId of the block having index &lt;x, y&gt; is (x + y * D<sub>x</sub>).
** For a grid of dimensions &lt;D<sub>x</sub>, D<sub>y</sub>&gt;, the blockId of the block having index &lt;x, y&gt; is (x + y * D<sub>x</sub>).
Line 84: Line 87:
** Arithmetic intensity and parallelism are paramount!
** Arithmetic intensity and parallelism are paramount!
** Memory-bound kernels require sufficiently high ''occupancy'' (the ratio of concurrently-running warps to maximum possible concurrent warps (as applied, usually, to [[#Streaming Multiprocessor|SMs]])) to hide latency.
** Memory-bound kernels require sufficiently high ''occupancy'' (the ratio of concurrently-running warps to maximum possible concurrent warps (as applied, usually, to [[#Streaming Multiprocessor|SMs]])) to hide latency.
* No branch prediction or speculation (and thus also no pipeline flushes on mispredicted branches).
* No branch prediction or speculation. Full predication.
{| border="1"
{| border="1"
! Memory type
! Memory type
Line 162: Line 165:


===Compute Capabilities===
===Compute Capabilities===
The original public CUDA revision was 1.0, implemented on the NV50 chipset corresponding to the GeForce 8 series. Compute capability, formed of a non-negative major and minor revision number, can be queried on CUDA-capable cards. All revisions thus far have been backwards-compatible.
The original public CUDA revision was 1.0, implemented on the NV50 chipset corresponding to the GeForce 8 series. Compute capability, formed of a non-negative major and minor revision number, can be queried on CUDA-capable cards. All revisions thus far have been fowards-compatible, though recent CUDA toolkits will not generate code for CC1 or 2.
 
{| border="1" class="wikitable"
{| border="1" class="wikitable"
! Resource
! Resource
Line 172: Line 176:
! 2.1 SM
! 2.1 SM
! 3.0 SMX
! 3.0 SMX
! 3.5 SMX
! 7.0 SM
! 7.5 SM
|-
|-
|CUDA cores
|CUDA cores
Line 181: Line 188:
|48
|48
|192
|192
|192
|64/32<br/>64/8
|64/2<br/>64/8
|-
|-
|Warp schedulers
|Schedulers
|1
|1
|1
|1
Line 189: Line 199:
|2
|2
|2
|2
|4
|4
|4
|4
|4
|-
|-
Line 199: Line 212:
|2
|2
|2
|2
|2
|1
|1
|-
|-
|Threads
|Threads
|768
|768
|768
|768
|1024
|1K
|1024
|1K
|1536
|1536
|1536
|1536
|2048
|2K
|2K
|2K
|1K
|-
|-
|Warps
|Warps
Line 217: Line 236:
|48
|48
|64
|64
|64
|64
|32
|-
|-
|Blocks
|Blocks
Line 226: Line 248:
|8
|8
|16
|16
|16
|32
|16
|-
|32-bit regs
|8K
|8K
|16K
|16K
|32K
|32K
|64K
|64K
|64K
|64K
|-
|-
|32-bit registers
|Examples
|8192
|G80
|8192
|G9x
|16384
|GT21x
|16384
|GT200
|32768
|GF110
|32768
|GF10x
|65536
|GK104
|GK110
|GV100
|TU10x
|-
|-
}
|}
{| border="1"
{| border="1"
! Revision
! Revision
Line 288: Line 328:
* 32 SFUs per SMX, 32 TFUs per ROP
* 32 SFUs per SMX, 32 TFUs per ROP
* 4 warp schedulers per SMX, capable of issuing two instructions per clock
* 4 warp schedulers per SMX, capable of issuing two instructions per clock
* Double-precision instructions can be paired with non-DP
** Previously, double-precision instructions couldn't be paired with anything
* ''PTX 3.0''
* ''PTX 3.0''
** <tt>madc</tt> and <tt>mad.cc</tt> instructions
** <tt>madc</tt> and <tt>mad.cc</tt> instructions
Line 297: Line 339:
** DWARF debugging symbols support
** DWARF debugging symbols support


|-
| 3.5
|
* 255 registers per thread
* "CUDA Dynamic Parallelism", the ability to spawn threads from within device code
* ''PTX 3.1''
** A funnel shift instruction, <tt>shf</tt>
** Loading read-only global data through the non-coherent texture cache, <tt>ld.global.nc</tt>
** 64-bit atomic/reduction operators extended to {or, xor, and, integer min, integer max}
** Mipmap type support
** Indirect texture/surface support
** Extends generic addressing to include the const state space
|-
| 7.0
|
* ''PTX 6.3''
* Tensor cores
* Independent thread scheduling
|-
| 7.5
|
* ''PTX 6.4''
* Integer matrix multiplication in tensor cores
|-
|-
|}
|}
Line 324: Line 391:
==Building CUDA Apps==
==Building CUDA Apps==
===nvcc flags===
===nvcc flags===
* <tt>-ptax-options=-v</tt> displays per-thread register usage
Pass flags to <tt>ptxas</tt> via -X:
* <tt>-X -v</tt> displays per-thread register usage
* <tt>-X -abi=no</tt> disables the PTX ABI, saving registers but taking away your stack
* <tt>-dlcm={cg,cs,ca}</tt> modifies cache behavior for loads
* <tt>-dscm={cw,cs}</tt> modifies cache behavior for stores
===SDK's common.mk===
===SDK's common.mk===
This assumes use of the SDK's common.mk, as recommended by the documentation.
This assumes use of the SDK's common.mk, as recommended by the documentation.
Line 376: Line 447:
* A run time limit is activated by default if the device is being used to drive a display.
* A run time limit is activated by default if the device is being used to drive a display.
* Please feel free to [mailto:nickblack@acm.org send me output!]
* Please feel free to [mailto:nickblack@acm.org send me output!]
{| border="1"
{| border="1"
! Device name
! Device name
Line 381: Line 454:
! MP's
! MP's
! Cores
! Cores
! Const mem
! Shmem/block
! Shmem/block
! Reg/block
! Reg/block
! Warp size
! Warp size
! Thr/block
! Thr/block
! Max pitch
! Texalign
! Texalign
! Clock
! Clock
Line 393: Line 464:
! Shared maps?
! Shared maps?
|-
|-
! COLSPAN="15" style="background:#8070D8;" | Compute capability 3.0
! COLSPAN="13" style="background:#eebeb6;" | Compute capability 7.0
|-
| Tesla V100
| 16GB
| 84
| 5376/2688/672
|
|
|
|
|
| 1.53GHz
| Yes
| No
| Yes
|-
! COLSPAN="13" style="background:#8070D8;" | Compute capability 3.0
|-
|-
| GeForce GTX 680
| GeForce GTX 680
Line 400: Line 487:
| 1536
| 1536
|  
|  
|
| 32
| 1024
|
|
|
|
Line 408: Line 492:
|
|
|
|
| Yes
| No
| No
| Yes
| Yes
|-
|-
! COLSPAN="15" style="background:#ffdead;" | Compute capability 2.1
! COLSPAN="13" style="background:#ffdead;" | Compute capability 2.1
|-
| GeForce GTX 560 Ti
|
|
|
|
|
|
|
|
|
|
|
|
|-
| GeForce GTX 550 Ti
|
|
|
|
|
|
|
|
|
|
|
|
|-
|-
| GeForce GTX 460
| GeForce GTX 460
Line 417: Line 530:
| 7
| 7
| 224
| 224
| 64k
| 48k
| 48k
| 32k
| 32k
| 32
| 32
| 1024
| 1024
| 2G
| 512b
| 512b
| 1.35GHz
| 1.35GHz
| 2x
| Yes
| No
| No
| Yes
| Yes
|-
|-
! COLSPAN="15" style="background:#ffdead;" | Compute capability 2.0
| GeForce GTS 450
|
|
|
|
|
|
|
|
|
|
|
|
|-
! COLSPAN="13" style="background:#ffdead;" | Compute capability 2.0
|-
|-
| GeForce GTX 580
| GeForce GTX 580
Line 435: Line 560:
| 16
| 16
| 512
| 512
| 1536
|
|
|
|
|
Line 442: Line 565:
| 1024
| 1024
|
|
|
| 1.544GHz
|
| Yes
| No
| No
| Yes
| Yes
Line 451: Line 574:
| 14
| 14
| 448
| 448
| 64k
| 48k
| 48k
| 32k
| 32k
| 32
| 32
| 1024
| 1024
| 2G
| 512b
| 512b
| 1.15GHz
| 1.15GHz
| 2x
| Yes
| No
| No
| Yes
| Yes
Line 467: Line 588:
| 14
| 14
| 448
| 448
| 64k
| 48k
| 48k
| 32k
| 32k
| 32
| 32
| 1024
| 1024
| 2G
| 512b
| 512b
| 1.15GHz
| 1.15GHz
| 2x
| Yes
| No
| No
| Yes
| Yes
Line 483: Line 602:
| 15
| 15
| 480
| 480
|
|
|  
|  
|  
|  
Line 499: Line 616:
| 14
| 14
| 448
| 448
|
|
|  
|  
|  
|  
Line 511: Line 626:
|
|
|-
|-
! COLSPAN="15" style="background:#efefef;" | Compute capability 1.3
! COLSPAN="13" style="background:#efefef;" | Compute capability 1.3
|-
|-
| Tesla C1060
| Tesla C1060
Line 517: Line 632:
| 30
| 30
| 240
| 240
| 65536b
| 16384b
| 16384b
| 16384
| 16384
| 32
| 32
| 512
| 512
| 262144b
| 256b
| 256b
| 1.30GHz
| 1.30GHz
Line 533: Line 646:
| 30
| 30
| 240
| 240
| 65536b
| 16384b
| 16384b
| 16384
| 16384
| 32
| 32
| 512
| 512
| 262144b
| 256b
| 256b
| 1.24GHz
| 1.24GHz
Line 549: Line 660:
| 30
| 30
| 240
| 240
| 65536b
| 16384b
| 16384b
| 16384
| 16384
| 32
| 32
| 512
| 512
| 262144b
| 256b
| 256b
| 1.48GHz
| 1.48GHz
Line 565: Line 674:
| 30
| 30
| 240
| 240
| 65536b
| 16384b
| 16384b
| 16384
| 16384
| 32
| 32
| 512
| 512
| 262144b
| 256b
| 256b
| 1.30GHz
| 1.30GHz
Line 581: Line 688:
| 27
| 27
| 216
| 216
| 65536b
| 16384b
| 16384b
| 16384
| 16384
| 32
| 32
| 512
| 512
| 262144b
| 256b
| 256b
| 1.47GHz
| 1.47GHz
Line 593: Line 698:
| Yes
| Yes
|-
|-
! COLSPAN="15" style="background:#efefef;" | Compute capability 1.2
! COLSPAN="13" style="background:#efefef;" | Compute capability 1.2
|-
|-
| GeForce GT 360M
| GeForce GT 360M
Line 599: Line 704:
| 12
| 12
| 96
| 96
| 65536b
| 16384b
| 16384b
| 16384
| 16384
| 32
| 32
| 512
| 512
| 262144b
| 256b
| 256b
| 1.32GHz
| 1.32GHz
Line 615: Line 718:
| 2
| 2
| 16
| 16
| 65536b
| 16384b
| 16384b
| 16384
| 16384
| 32
| 32
| 512
| 512
| 262144b
| 256b
| 256b
| 1.40GHz
| 1.40GHz
Line 631: Line 732:
| 12
| 12
| 96
| 96
| 65536b
| 16384b
| 16384b
| 16384
| 16384
| 32
| 32
| 512
| 512
| 262144b
| 256b
| 256b
| 1.424GHz
| 1.424GHz
Line 643: Line 742:
| Yes
| Yes
|-
|-
! COLSPAN="15" style="background:#efefef;" | Compute capability 1.1
! COLSPAN="13" style="background:#efefef;" | Compute capability 1.1
|-
|-
| ION
| ION
Line 649: Line 748:
| 2
| 2
| 16
| 16
| 65536b
| 16384b
| 16384b
| 8192
| 8192
| 32
| 32
| 512
| 512
| 262144b
| 256b
| 256b
| 1.1GHz
| 1.1GHz
Line 665: Line 762:
| 2
| 2
| 16
| 16
| 65536b
| 16384b
| 16384b
| 8192
| 8192
| 32
| 32
| 512
| 512
| 262144b
| 256b
| 256b
| 0.92GHz
| 0.92GHz
Line 681: Line 776:
| 16
| 16
| 128
| 128
| 65536b
| 16384b
| 16384b
| 8192
| 8192
| 32
| 32
| 512
| 512
| 2147483647b (!)
| 256b
| 256b
| 1.84GHz
| 1.84GHz
Line 697: Line 790:
| 16
| 16
| 128
| 128
| 65536b
| 16384b
| 16384b
| 8192
| 8192
| 32
| 32
| 512
| 512
| 262144b
| 256b
| 256b
| 1.67GHz
| 1.67GHz
Line 713: Line 804:
| 8
| 8
| 64
| 64
| 65536b
| 16384b
| 16384b
| 8192
| 8192
| 32
| 32
| 512
| 512
| 262144b
| 256b
| 256b
| 1.62GHz,
| 1.62GHz,
Line 730: Line 819:
| 2
| 2
| 16
| 16
| 65536b
| 16384b
| 16384b
| 8192
| 8192
| 32
| 32
| 512
| 512
| 262144b
| 256b
| 256b
| 0.88GHz
| 0.88GHz
Line 746: Line 833:
| 16
| 16
| 128
| 128
| 65536b
| 16384b
| 16384b
| 8192
| 8192
| 32
| 32
| 512
| 512
| 262144b
| 256b
| 256b
| 1.62GHz
| 1.62GHz
Line 762: Line 847:
| 4
| 4
| 32
| 32
| 65536b
| 16384b
| 16384b
| 8192
| 8192
| 32
| 32
| 512
| 512
| 262144b
| 256b
| 256b
| 0.95GHz
| 0.95GHz
Line 778: Line 861:
| 1
| 1
| 8
| 8
| 65536b
| 16384b
| 16384b
| 8192
| 8192
| 32
| 32
| 512
| 512
| 262144b
| 256b
| 256b
| 1.40GHz
| 1.40GHz
Retrieved from "https://nick-black.com/dankwiki/index.php?title=CUDA"