对cuda内核的一点认识

最新推荐文章于 2024-03-29 18:05:15 发布

Zino-齐诺

最新推荐文章于 2024-03-29 18:05:15 发布

阅读量1k

点赞数

文章标签： cuda

cuda的基本运行原理参考：https://www.cnblogs.com/skyfsm/p/9673960.html

关于cuda的内存问题

在GPU上CUDA线程可以访问到的存储资源有很多，每个CUDA线程拥有独立的本地内存（local Memory）；每一个线程块（block）都有其独立的共享内存（shared memory），共享内存对于线程块中的每个线程都是可见的，它与线程块具有相同的生存时间；同时，还有一片称为**全局内存（global memory）**的区域对所有的CUDA线程都是可访问的。

除了上述三种存储资源以外，CUDA还提供了两种只读内存空间：常量内存（constant memory）和纹理内存（texture memory），同全局内存类似，所有的CUDA线程都可以访问它们。对于一些特殊格式的数据，纹理内存提供多种寻址模式以及数据过滤方法来操作内存。这两类存储资源主要用于一些特殊的内存使用场合。

一个程序启动内核函数以后，全局内存、常量内存以及纹理内存将会一直存在直到该程序结束。

作为cpu的一个协处理器，为了完成和cpu的信息传递，CUDA程序管理两大块由DRAM构成的内存区域：CPU端可以访问到的主机内存（host memory）以及GPU端供CUDA内核访问到的设备内存（device memory），设备内存主要由全局内存、常量内存以及纹理内存构成。

关于cuda编译的问题

离线编译部分：在使用cmake时，在CMakelist.txt中添加

find_package(CUDA REQUIRED)
include_directories( ${CUDA_INCLUDE_DIRS})

对于包含了cuda代码的文件在链接时，要将

add_library(...)

改为

cuda_add_library(...)

另外，如要使用c++11编译含有cuda的代码，在开头添加

list(APPEND CUDA_NVCC_FLAGS "-std=c++11")

补充部分：
即时编译（Just-in-time Compile）部分会在程序装载时自动完成，该部分会曾在程序装载的时间，目的是使得编译好的程序可以从新的显卡驱动中获得性能提升。同时到目前为止，这一方法是保证编译好的程序在还未问世的GPU上运行的唯一解决方案。
在即时编译的过程中，显卡驱动将会自动缓存PTX代码的编译结果，以避免多次调用同一程序带来的重复编译开销。NVIDIA把这部分缓存称作计算缓存（compute cache），当显卡驱动升级时，这部分缓存将会自动清空，以使得程序能够自动获得新驱动为即时编译过程带来的性能提升。
在即时编译的过程中，显卡驱动将会自动缓存PTX代码的编译结果，以避免多次调用同一程序带来的重复编译开销。NVIDIA把这部分缓存称作计算缓存（compute cache），当显卡驱动升级时，这部分缓存将会自动清空，以使得程序能够自动获得新驱动为即时编译过程带来的性能提升。

有一些环境变量可以用来控制即时编译过程：

设置CUDA_CACHE_DISABLE为1将会关闭缓存功能
CUDA_CACHE_MAXSIZE变量用于指定计算缓存的字节大小，默认情况下它的值是32MB，它最大可以被设置为4GB。任何大于缓存最大值得二进制代码将不会被缓存。在需要的情况下，一些旧的二进制代码可能被丢弃以腾出空间缓存新的二进制代码。
CUDA_CACHE_PATH变量用于指定计算缓存的存储目录地址，它的缺省值如下：
设置CUDA_FORCE_PTX_JIT为1会强制显卡驱动忽略应用程序中的二进制代码并且即时编译程序中的嵌入PTX代码。如果一个内核函数没有嵌入的PTX代码，那么它将会装载失败。该变量可以用来确认程序中存在嵌入的PTX代码。同时，使用即时编译（just-in-time Compilation）技术也可确保程序的向前兼容性。

如何设置grid和block的参数

通常我们在调用kernel核函数的时候，总的核函数个数是固定的，即gridNum × blockNum，我们通过调整gridNum和blockNum的大小来使gpu上的硬件资源被尽可能的发挥出来。
gpu硬件结构
同一个block里的所有线程都会被分配到一个sm里，以CUDA Fermi硬件为例，一个sm里有8个sp，因此一个sm里最多被分配8个block。每一个block中的线程以32个为一组组成一个warp，warp之间的调度没有额外开销。
注意：warp是调度和执行的基本单位，half-warp是存储器操作的基本单位。
二维的demo

// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N],
float C[N][N])
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    if (i < N && j < N)
        C[i][j] = A[i][j] + B[i][j];
}

int main()
{
    ...
    // Kernel invocation
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}

补充：
1.当要分配二维或三维内存空间时可以使用cudaMallocPitch()和cudaMalloc3D()函数来分配线性内存。这些函数能够确保分配的内存满足设备内存访问的对齐要求，对于行地址的访问以及多维数组间的数据传输提供高性能保证。

参考：
1、https://docs.nvidia.com/cuda/
2、https://blog.csdn.net/curtern/article/details/81273024
3、https://www.cnblogs.com/biglucky/p/4346554.html
4、https://blog.csdn.net/qq_41598072/article/details/81086500
5、https://blog.csdn.net/dcrmg/article/details/54867507

小觅gpu参数：https://www.nvidia.cn/object/tegra-superchip-cn.html