本地,全局,常量和共享内存

Dou*_*oug 2 cuda gpu-shared-memory gpu-constant-memory gpu-local-memory

我读了一些引用本地内存的CUDA文档.(这主要是早期文档.)设备属性报告本地mem大小(每个线程)."本地"记忆是什么意思?什么是'本地'记忆?"本地"记忆在哪里?如何访问"本地"内存?这是__device__记忆,不是吗?

设备属性还报告:全局,共享和常量mem大小.这些陈述是否正确: 全局记忆是__device__记忆.它具有网格范围和网格(内核)的生命周期. 常量内存是__device__ __constant__内存.它具有网格范围和网格(内核)的生命周期. 共享内存是__device__ __shared__内存.它具有单个块范围和该块(线程)的生命周期.

我在想共享内存是SM内存.即只有那个SM才能直接访问的内存.资源相当有限.SM是不是一次分配了一堆块?这是否意味着SM可以交错执行不同的块(或不是)?即运行阻止*A*线程直到它们停止.然后运行block*B*线程直到它们停止.然后再次换回阻止*A*线程.或者SM是否为块*A*运行一组线程,直到它们停止.然后交换另一组块*A*线程.此交换继续,直到块*A*用尽.然后,只有这样才能在块*B*上开始工作.因为共享记忆,我问.如果单个SM从2个不同的块交换代码,那么SM如何快速交换/分出共享内存块?(我认为后来的senerio是真的,并且没有交换共享内存空间.块*A*运行直到完成,然后块*B*开始执行.注意:块*A*可能是不同的内核比块*B*.)

lim*_*mes 5

从CUDA C编程指南第5.3.2.2节,我们看到在几种情况下使用本地内存:

  • 当每个线程都有一些数组,但它们的大小在编译时是不知道的(因此它们可能不适合寄存器)
  • 在编译时已知数组的大小,并且此大小对于寄存器内存来说太大(这也可能发生在大结构上)
  • 当内核已经使用了所有的寄存器内存(所以如果我们充满了寄存器n int秒,n+1int会进入本地内存) -这最后一种情况是寄存器溢出,并且应该避免,因为:

"本地"存储器实际上存在于全局存储空间中,这意味着与寄存器和共享存储器相比,对它的读取和写入相对较慢.每次在内核中使用一些不适合寄存器的变量,数组等时,您将访问本地内存,不是共享内存,并且不作为全局内存传递.你不必做任何明确的事情来使用它 - 实际上你应该尽量减少它的使用,因为寄存器和共享内存要快得多.

编辑:Re:共享内存,你不能有两个块交换共享内存或查看彼此的共享内存.由于无法保证块的执行顺序,如果您尝试这样做,可能需要将SMP占用几个小时,等待另一个块执行.同样,同时在设备上运行的两个内核无法看到彼此的内存,除非它是全局内存,即便如此,你也在玩火(竞争条件).据我所知,块/内核无法真正向对方发送"消息".你的场景并没有真正意义,因为块的执行顺序每次都不同,停止一个块等待另一个块是不好的做法.