cuda理解

RTX3050GPU显卡详细信息

Detected 1 CUDA Capable device(s)
Device 0: "NVIDIA GeForce RTX 3050"
 CUDA Driver Version / Runtime Version 12.1 / 11.6
 CUDA Capability Major/Minor version number: 8.6
 Total amount of global memory: 8.00 MBytes (8589410304 bytes)
 GPU Clock rate: 1777 MHz (1.78 GHz)
 Memory Clock rate: 7001 Mhz
 Memory Bus Width: 128-bit
 L2 Cache Size: 1572864 bytes   //1.5M
 Max Texture Dimension Size (x,y,z)  1D=(131072), 
                                     2D=(131072,65536),
                                     3D=(16384,16384,16384)
 Max Layered Texture Size (dim) x layers 1D = ( 32768) x  2048,
                                         2D = ( 32768,  32768) x  2048
 Total amount of constant memory: 65536 bytes           //64kB
 Total amount of shared memory per block: 49152 bytes   //48kB
 Total number of registers available per block: 65536   // 寄存器数量
 Warp size: 32   // 一个warp 由32个线程组成，1536/32=48
 Maximum number of threads per multiprocessor: 1536     // 每个SM单元中最多有1536个线程
 Maximum number of threads per block: 1024  // 每个block包含的的线程数为1024个
 Maximum sizes of each dimension of a block: 1024 x 1024 x 64   // 每个block中各个维度最大设置的维度，
                                                                // example:1024*1*1, 1*1024*1, 4*4*64
                                                                // 总数计算不能超过单个block的最大数
 Maximum sizes of each dimension of a grid: 2147483647 x 65535 x 65535  // 
 Maximum memory pitch: 2147483647 bytes
Vector size 32
Execution configuration<<<1,1>>>
Array matched.

cudaMalloc的理解

cudaMalloc的原型为

cudaError_t cudaMalloc(void **devPtr, size_t size)

这个原型与C语言中的malloc大同小异，

int *a = (int *)malloc(n*sizeof(int))，返回的是一个int型指针，指向大小为n个int型数据的连续内存地址的首地址，可以理解为a是这个数组的首地址。

在CPU中通常使用malloc函数分配内存，命令如下所示，

float* tmp = (float*)malloc(sizeof(float));
*tmp = 10;
printf("%f",tmp);  // 输出tmp所在的首地址
printf("%f",*tmp); // 输出tmp里面的值

我们在进行cuda编程的时候，第一步需要在GPU内分配内存，与数组的声明的步骤是一样的。假如我们要在GPU内申明端n个大小的float数组，我们需要定义float *addr，用于指向GPU内这个地址的首地址。因此，addr这个变量中存的就是用户在GPU中声明的float型数组的首地址。

float *addr;
int nBytes = 1024;
cudaMalloc((float**)&addr, nBytes);

第二个参数大家都知道具体意思了，表示数据的大小。

第一个参数我们首先看到的是&addr。前面我们假设addr内部存放的是一个指向GPU中的数组的首地址，需要对addr的内容进行改变，因此我们必须采用引用的方式进行形参传递。否则，完成了cudaMalloc函数调用后仅仅是形参的值发生变化（参考C语言学习中的引用调用），原来我们声明的float *addr这个实参没有发生任何变化。注意：我们需要的是把addr指针的内容改变，而函数原型的devPtr返回的是地址的值，因此需要通过对addr这个指针进行引用调用，改变其内部的值。

(float**)&addr// flaot(**)+(&addr)==** devPtr

理解了上面的内容，接着是强制类型转换成指针型指针的理解，就简单多了。前面说过，addr是指向地址的指针。cudaMalloc完成了*addr的内容的改变后，需要转换数据类型。把它转换成指针型指针是对于主机端而言的（GPU称为设备端），addr这个变量是指向我们在GPU内部声明的连续地址的首地址，因此，我们对addr进行第一次引用计算，得到的是首地址的值。我们需要通过这个值来在GPU的内存进行操作，因此需要再做一次引用计算，得到的就是GPU中连续地址的第一个单元，接下来就可以进行主机端设备端的内存内容拷贝了。

总结
cudaMalloc的参数与我们设置的地址类型数据有一定的关系。由于cudaMemcpy等重要的内存操作函数均以指针作为形参，因此我们需要定义一个与之相同类型的指针完成形参赋值。

由于指针是指向地址的，必须把我们声明的指针变量初始化为GPU的内存首地址。这样一来，addr实质上就变成了一个存放了地址的指针。这也决定了cudaMalloc函数调用时，需要对addr进行引用计算，将GPU内存首地址放在addr指向的地址中。完成了这个函数调用后，addr就变成了一个指向GPU内存首地址的地址（指针）了。