2.目前一个kernel中只有一个grid,grid只能设定x、y维,z维默认为1,但是block的三维都可以自己设定。
3.一个block中最多有512个thread,这些thread会被分割为线程束warp在sm上进行调度。每个warp包含连续的32个线程,访问存储器时是以半线程束为单位的。
4.cuda的源文件必须使用nvcc编译器进行编译。这个编译器在编译的时候,会分离出主机端和设备端代码,主机端调用其他高性能编译器编译,如gcc,设备端代码由nvcc编译成ptx代码或者二进制代码。ptx代码是为动态编译器JIT设计的,这样可以应付不同显卡上的不同的机器语言。
5.在计算能力1.3的设备中,每个sm的寄存器文件数量为16384,每个寄存器文件大小为32bit,如果每个线程使用的私有变量太多或者大小不定,可能将这样的变量分配到局部寄存器中(local memory)。例如usigned int mt[3]就是局部寄存器的。 变量被分在哪里可以通过编译输出的ptx汇编代码查看,有.local的则为局部寄存器中的。
6.共享存储器:extren __shared__声明的变量都开始于同一个地址,因此变量的布局必须通过显式的定义,如
extern __shared__ char array[];
__device__ void func()
{
short* array0 = (short*)array;
float* array1 = (float*)&array0[128];
int array2 = (int*)&array1[64];
}
7. cuda主机端的内存被分为两种:可分页内存(pageable memory)(由malloc和new分配)和页锁定(pinned)内存(一定是存储在物理内存,不会存储在虚拟内存,因此速度较快,可以通过DMA加速与设备通信,使用cudaHostAlloc 和 cudaFreeHost分配和释放
)
cuda2.3版本以上支持portable memory,在
cudaHostAlloc()时加上cudaHostAllocPortable标志,可以使多个CPu共享一块页锁定内存以实现CPU间的通信。默认情况下,pinned内存只能由分配的线程访问。
write-combined Memory,在分配
cudaHostAlloc()时设定cudahostAllocWriteCombin ed的标志,可以免除cpu对内存的监视,减少将内存上的数据缓存到L1、L2 cache中,在主机和显存数据传输时节省了时间,这样的作法比较适合于主机端只写的数据。
mapped memory,主机端指针通过
cudaHostAlloc()获得,设备端通过cudaHostGetDevicePointer ()获得,这样可以从kernel里直接访问主存,省略了在显存上分配空间,实现了对内存的零拷贝访问,不过使用前要通过cudaGetDeviceProperties()判断设备是否有canMapHostMemory的属性。
8.常数存储:显存上的只读地址空间,拥有缓存加速,每个sm拥有8KB的常数存储器缓存。
__constant__ char p_hello[];
可以通过cudaMemcpyToSymbol(p_hello, hello_host, sizeof(char)*11)从内存把数据拷进去。
9.纹理存储器:只读存储器,具备一些特殊的功能。
<script type="text/javascript" id="wumiiRelatedItems"> </script>