CUDA编程学习

六树蚊

已于 2024-08-22 01:32:43 修改

阅读量248

点赞数 2

文章标签：学习人工智能

于 2024-06-17 21:50:07 首次发布

本文链接：https://blog.csdn.net/weixin_55533898/article/details/139754724

版权

1. 函数类型限定符

__global__
表示一个内核函数，主机（CPU）上调用，但在设备（GPU）上执行，同定义的多个cuda线程同时执行，在<<<>>>中设置cuda线程数，每个线程有专属id

__host__
由CPU调用，且在CPU执行的函数，默认的函数限定符

__device__
设备执行的程序，device前缀定义的函数只能在GPU上执行，所以device修饰的函数里面不能调用一般常见的函数

2. 变量类型限定符
device

__shared__

__constant__

3. cuda线程概念

参考资料：CUDA学习笔记_cuda dim3-CSDN博客

进程是系统资源分配的基本单位，而线程是CPU 能调度和独立运行的基本单位（最小单元）

3.1 Grid、Block和Thread

CUDA的软件架构由网格（Grid）、线程块（Block）和线程（Thread）组成，分解出的多线程需要能够独立并行运行，不相互依赖

每个block最多包含1024个thread，共享处理器的内存资源

block可以组织成一维、二维或三维的线程块网格(grid)

三者的关系如下图：

3.2 <<<...>>>

<<<block, thread>>> 类型为int或者dim3，定义cuda的多线程数量，dim3可以指定二维，可以通过内置blodkId获得线程索引

3.2.1 dim3数据结构

用dim3指定二维

dim3 threadsPerBlock(16, 16);
dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);

线程块大小为16x16(256个线程)

constexpr int BLOCK_SIZE = 256;
dim3 blocks = ptnums/256 + 1;
dim3 threads = BLOCK_SIZE;
checkCudaErrors(cudaStreamSynchronize(cudaStreamPerThread));
Function<<<blocks, threads>>>(lidarPt);
checkCudaErrors(cudaStreamSynchronize(cudaStreamPerThread));
checkCudaErrors(cudaGetLastError());

-- Function为函数

-- <<<blocks, threads>>>表示调用的blocks和threads的数量

-- lidarPt为输入的参数

__global__ void PreprocPc3l(pt)
{
    const int threadId = blockIdx.x * blockDim.x + threadIdx.x;

    if (threadId >= pt.size()) {
        return;
    }
}

CUDA中每一个线程都有一个唯一的标识ID—ThreadIdx，这个ID随着Grid和Block的划分方式的不同而变化

线程同步

Cuda Streams的概述（四）-- 同步_cudastreamsynchronize-CSDN博客

void test()
{
    constexpr int BLOCK_SIZE = 256;
    dim3 stackBlockPerGrid = DIVE_UP_3L(KVoxelMaxArea3l, BLOCK_SIZE);
    fun1<<<stackBlockPerGrid, BLOCK_SIZE>>>(input);
    fun2<<<stackBlockPerGrid, BLOCK_SIZE>>>(input);
    CheckOdomGradientValid3l<<<stackBlockPerGrid, BLOCK_SIZE>>>(input);
    checkCudaErrors3l(cudaGetLastError());
//    checkCudaErrors3l(cudaStreamSynchronize(cudaStreamPerThread));
    checkCudaErrors3l(cudaDeviceSynchronize());
}

这边应该用checkCudaErrors3l(cudaDeviceSynchronize());，用checkCudaErrors3l(cudaStreamSynchronize(cudaStreamPerThread));这边cuda并行的现成还没走完，就会走test之后的代码，会导致拷贝计算赋值等操作不完全（巨心累的一天.......）