cuda中local memory的上限问题
CUDA中local memory的上限问题
前言
在CUDA编程的过程中，核函数以及device属性的函数内声明的局部变量会有限存储在线程独有的寄存器上，若寄存器容纳不下，比如申请了 一个较大的数组，则会存储在local memory中，在一些书籍和文档中有提到过，local memory是一种特殊的global memory，当local memory不足时，则申请的变量或数组会被存储在global memory中，但我在实际编程过程中却发现溢出的部分并不会被存储在global memory中，而是离奇地消失了，而且这种情况下编译和运行甚至都不会报错，只会因为该变量没有值而导致运算结果错误。目前我尚不清楚导致此现象的机制，但由此可以肯定的是local memory是存在上限的，且存储上限并不大。
local memory的上限测试
测试环境
官方文档中给出的local memory的上限为512KB,但实际测试发现并不能达到512KB，测试环境为：
- 设备： NVIDIA A30
- SM数： 56
- Max threads per SM： 2048
- 显存： 24258M
CUDA C
CUDA C测试代码：
#include<stdio.h>
__global__ void localmemtest()
{
float a[53][1024];
int i,j;
for(i=0;i<53;i++)
for(j=0;j<1024;j++)
a[i][j]=blockIdx.x+threadIdx.x;
if(threadIdx.x == 0 && threadIdx.y == 0)
printf("%f\n",a[1][1]);
while(1)
{}
}
int main()
{
dim3 dimgrid(1,1,1);
dim3 dimblock(32,1,1);
localmemtest<<<dimgrid,dimblock>>>();
cudaThreadSynchronize();
return 0;
}
编译方式为:
nvcc -O0 --threads 0 -gencode arch=compute_80,code=sm_80 test.cu
实际测试过程中发现，即使加了-O0选项，如果不对每个元素进行赋值或做一些其它操作，编译器仍会自动进行优化而导致测试结果不准确，因此增加了循环赋值语句避免编译器优化，另外增加一个死循环便于观察显存占用情况。编译选项中--threads 0代表开启一个线程0进行编译，-gencode arch=compute_80,code=sm_80作用为指明显卡的计算能力（compute capbility），生成相应的机器码，A30的计算能力为8.0，若计算能力不为8.0，则需要调整这里出现的两个数字进行适配，另外，CUDA C编译过程中如果不加此编译选项是无法生成正确的可运行的机器码的。
测试结果
经测试发现，此二维数组最大只能达到53×1024个float或者int元素，也即212KB左右，改用double类型，测得结果近似。若要更精确地测量应该使用字节数更小的char类型，但由于CUDA 对char类型的支持很弱，不好分辨其空间是否正常分配，因此可以选择short或int等进行测试。short应该更加精确。
CUDA Fortran
换用CUDA Fortran进行测试，测试代码为：
attributes(global) subroutine localmemtest()
implicit none
real,dimension(53,1024) :: a
integer :: i,j
do i = 1, 1024
do j = 1, 53
a(j,i) = blockidx%x + threadidx%x
end do
end do
if(threadidx%x .eq. 1 .and. threadidx%y .eq. 1) then
print *,a(1,1)
end if
end subroutine localmemtest
program main
use cudafor
implicit none
type(dim3) :: dimgrid,dimblock
integer ierr
dimgrid = dim3(1,1,1)
dimblock = dim3(1,1,1)
call localmemtest<<<dimgrid,dimblock>>>()
ierr = cudathreadsynchronize()
call sleep(100)
end program main
编译方式为：
nvfortran -O0 test.cuf
CUDA Fortran不需要指明计算能力，直接编译即可。
测试结果
CUDA Fortran的测试结果和C是一致的，均为最多申请53K（K=1024）个4字节数据类型的空间。总大小约为212KB。
官方文档参见：
关于local memory limit的图表在section G.1 table 12
local memory的上限
可以看到我们测试所得的结果和官方文档上给出的并不一致，通过查阅资料找到了local memory limit的计算公式为：
min(amount of loacl memory per thread as documented in seciton G.1 table 12, available GPU memory / number of SMs / maximum resident threads per SM )
由我在前面给出的测试环境数据，代入到此公式内进行计算：
available GPU memory / number of SMs / maximum resident threads per SM = 24258MB / 56 / 2048 = 0.2115MB = 211.5KB
和我们测试所得结果一致。其他版本的GPU也可代入此公式计算其local memory的上限大小。
