CUDA编程基本使用

最新推荐文章于 2026-03-29 10:34:01 发布

原创最新推荐文章于 2026-03-29 10:34:01 发布 · 1.5k 阅读

28 ·

本内容遵循CC 4.0 BY-SA版权协议

标签

#python

AudioSeal 音频水印系统

**AudioSeal** 是 Meta 开源的语音水印系统，用于 AI 生成音频的检测和溯源。

文章目录

获取设备相关信息
helloWorld
线程和块的索引
并行计算

前期回顾：
CUDA编程入门介绍

获取设备相关信息

在CUDA编程中获取设备相关信息至关重要，体现在设备兼容性和可用性检查以及性能优化和资源合理利用两方面，包括确保CUDA支持、确定设备ID以获取属性、了解硬件能力、掌握内存信息和利用缓存及纹理特性等。

官网有相关属性的解释：[https://docs.nvidia.com/cuda/cuda-runtime-api/structcudaDeviceProp.html#structcudaDeviceProp]

int main(int argc,char **argv){
    int deviceCount=0;
    cudaGetDeviceCount(&deviceCount);//获取当前可用的CUDA设备数量

    if (deviceCount == 0){
        printf("There are no available device(s) that support CUDA\n");
    }
    else{
        printf("Detected %d CUDA Capable device(s)\n", deviceCount);
    }
    //dev表示要查询的设备ID（设为0），driverVersion和runtimeVersion用于存储CUDA驱动程序版本和运行时版本。
    int dev = 0, driverVersion = 0, runtimeVersion = 0;
    cudaSetDevice(dev);//设置要使用的CUDA设备为dev（设备0）
    cudaDeviceProp deviceProp;//声明cudaDeviceProp结构体变量deviceProp，用于存储设备的属性。
    cudaGetDeviceProperties(&deviceProp, dev);//获取设备的详细属性，并将其存储在deviceProp中。
    printf("Device %d: \"%s\"\n", dev, deviceProp.name);//输出设备编号和设备名称

    cudaDriverGetVersion(&driverVersion);//获取CUDA驱动程序版本
    cudaRuntimeGetVersion(&runtimeVersion);//获取CUDA运行时版本
    
    printf("  CUDA Driver Version / Runtime Version          %d.%d / %d.%d\n",
           driverVersion / 1000, (driverVersion % 100) / 10,
           runtimeVersion / 1000, (runtimeVersion % 100) / 10);
    //输出设备的CUDA（Compute Capability）主版本号和次版本号。
    printf("  CUDA Capability Major/Minor version number:    %d.%d\n",deviceProp.major, deviceProp.minor);
    //设备的全局内存总量（以GB和字节为单位）
    printf("  Total amount of global memory:                 %.2f GBytes (%llu bytes)\n", 
            (float)deviceProp.totalGlobalMem / pow(1024.0, 3),
            (unsigned long long)deviceProp.totalGlobalMem);
    //GPU时钟频率（以MHz和GHz为单位）
    printf("  GPU Clock rate:                                %.0f MHz (%0.2f GHz)\n", deviceProp.clockRate * 1e-3f,deviceProp.clockRate * 1e-6f);
    //设备的内存时钟频率（以MHz为单位）和内存总线宽度（以位为单位）
    printf("  Memory Clock rate:                             %.0f Mhz\n",deviceProp.memoryClockRate * 1e-3f);
    printf("  Memory Bus Width:                              %d-bit\n",deviceProp.memoryBusWidth);
    //如果设备有L2缓存，打印L2缓存的大小
    if (deviceProp.l2CacheSize){
        printf("  L2 Cache Size:                                 %d bytes\n",deviceProp.l2CacheSize);
    }
    //设备支持的最大纹理尺寸：1D、2D和3D
    printf("  Max Texture Dimension Size (x,y,z)             1D=(%d), 2D=(%d,%d), 3D=(%d,%d,%d)\n", 
            deviceProp.maxTexture1D,
            deviceProp.maxTexture2D[0], deviceProp.maxTexture2D[1],
            deviceProp.maxTexture3D[0], deviceProp.maxTexture3D[1],
            deviceProp.maxTexture3D[2]);
    //层纹理的最大尺寸和层数
    printf("  Max Layered Texture Size (dim) x layers        1D=(%d) x %d, 2D=(%d,%d) x %d\n", 
            deviceProp.maxTexture1DLayered[0],
            deviceProp.maxTexture1DLayered[1], deviceProp.maxTexture2DLayered[0],
            deviceProp.maxTexture2DLayered[1],
            deviceProp.maxTexture2DLayered[2]);
    //常量内存总量、每个块的共享内存、每个块的可用寄存器数量和warp大小
    printf("  Total amount of constant memory:               %zu bytes\n",deviceProp.totalConstMem);
    printf("  Total amount of shared memory per block:       %zu bytes\n",deviceProp.sharedMemPerBlock);
    printf("  Total number of registers available per block: %d\n",deviceProp.regsPerBlock);
    printf("  Warp size:                                     %d\n",deviceProp.warpSize);
    //每个多处理器支持的最大线程数,共享内存和每个块的最大线程数
    printf("  registers available per multiprocessor:        %u bytes\n",deviceProp.regsPerMultiprocessor);
    printf("  Shared memory available per multiprocessor:    %zu bytes\n",deviceProp.sharedMemPerMultiprocessor);
    printf("  Maximum number of threads per multiprocessor:  %d\n",deviceProp.maxThreadsPerMultiProcessor);
    printf("  Maximum number of threads per block:           %d\n",deviceProp.maxThreadsPerBlock);
    //块和网格每个维度的最大尺寸
    printf("  Maximum sizes of each dimension of a block:    %d x %d x %d\n",
           deviceProp.maxThreadsDim[0],deviceProp.maxThreadsDim[1],deviceProp.maxThreadsDim[2]);
    printf("  Maximum sizes of each dimension of a grid:     %d x %d x %d\n",
           deviceProp.maxGridSize[0],deviceProp.maxGridSize[1],deviceProp.maxGridSize[2]);
    //设备的最大内存填充（pitch）大小
    printf("  Maximum memory pitch:                          %zu bytes\n",deviceProp.memPitch);
    
    exit(EXIT_SUCCESS);
}

helloWorld

并行执行的实现：定义helloFromGPU函数，用__global__标识，定义一个block，5个线程并行执行。

__global__：表示这是一个在GPU上运行的CUDA核函数，可以从主机端（CPU）调用。
cudaDeviceReset();：在程序结束之前，重置CUDA设备，释放资源。

#include<stdio.h>

__global__ void helloFromGPU(){
    printf("Hello from GPU,threadId:(%d,%d,%d)\n", threadIdx.x, threadIdx.y, threadIdx.z);
}

int main(){
    printf("Hello from CPU\n");
    helloFromGPU<<<1,5>>>();
    cudaDeviceReset();
    return 0;
}

输出：

Hello from CPU
Hello from GPU,threadId:(0,0,0)
Hello from GPU,threadId:(1,0,0)
Hello from GPU,threadId:(2,0,0)
Hello from GPU,threadId:(3,0,0)
Hello from GPU,threadId:(4,0,0)

可见执行了5个线程，并打印出了各个线程的id，block的维度是（5，1，1）

加入约束条件

约束前3个线程和后面的线程执行不同的任务。如下面的代码：首先获取全局线程ix，只需要用条件判断即刻实现不同线程执行不同的内容。

#include<stdio.h>

__global__ void helloFromGPU(){
    unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;//全局线程索引
    if (ix<3)
    printf("Hello from GPU, above threadId:%d\n", threadIdx.x);
    if (ix>=3)
    printf("Hello from GPU, under threadId:%d\n", threadIdx.x);
}

int main(){
    printf("Hello from CPU\n");
    helloFromGPU<<<1,5>>>();
    cudaDeviceReset();
    return 0;
}

输出：

Hello from CPU
Hello from GPU, above threadId:0
Hello from GPU, above threadId:1
Hello from GPU, above threadId:2
Hello from GPU, under threadId:3
Hello from GPU, under threadId:4

通过约束线程索引，上面GPU前三个线程和后两个不一样。

线程和块的索引

通过一个CUDA核函数（__global__）打印出每个线程所在的块的索引和线程在该块中的索引。

gridDim：内置变量，表示整个CUDA网格的尺寸，包含x、y和z三个维度。
blockDim：内置变量，表示每个块的尺寸，包含x、y和z三个维度。
blockIdx：内置变量，表示当前块在网格中的索引（块的索引），同样是三个维度x、y、z。
threadIdx：内置变量，表示当前线程在块中的索引（线程的索引），也是三维的。

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void checkIndex(void)
{
    // gird维度(block的排列)，block维度（thread的排列）
    //打印block的索引，和该block中的线程索引
    printf("gridDim:(%d, %d, %d) blockDim:(%d, %d, %d) blockIdx:(%d, %d, %d) threadIdx:(%d, %d, %d)\n", 
        gridDim.x, gridDim.y, gridDim.z, 
        blockDim.x, blockDim.y, blockDim.z, 
        blockIdx.x, blockIdx.y, blockIdx.z, 
        threadIdx.x, threadIdx.y, threadIdx.z);
}

int main(int argc, char **argv)
{
    // 数据大小
    int nElem = 6;

    // block排列和线程排列
    dim3 block(3);
    dim3 grid((nElem + block.x - 1) / block.x); //等同于向上取整，得到grid的大小

    // 从主机侧检查网格和块尺寸
    printf("grid.x %d grid.y %d grid.z %d\n", grid.x, grid.y, grid.z);
    printf("block.x %d block.y %d block.z %d\n", block.x, block.y, block.z);

    // 从设备侧检查网格和块尺寸
    checkIndex<<<grid, block>>>();
    cudaDeviceReset();

    return(0);
}

输出：

grid.x 2 grid.y 1 grid.z 1
block.x 3 block.y 1 block.z 1
gridDim:(2, 1, 1) blockDim:(3, 1, 1) blockIdx:(0, 0, 0) threadIdx:(0, 0, 0)
gridDim:(2, 1, 1) blockDim:(3, 1, 1) blockIdx:(0, 0, 0) threadIdx:(1, 0, 0)
gridDim:(2, 1, 1) blockDim:(3, 1, 1) blockIdx:(0, 0, 0) threadIdx:(2, 0, 0)
gridDim:(2, 1, 1) blockDim:(3, 1, 1) blockIdx:(1, 0, 0) threadIdx:(0, 0, 0)
gridDim:(2, 1, 1) blockDim:(3, 1, 1) blockIdx:(1, 0, 0) threadIdx:(1, 0, 0)
gridDim:(2, 1, 1) blockDim:(3, 1, 1) blockIdx:(1, 0, 0) threadIdx:(2, 0, 0)

可知grid维度为2（2个块），每个块的维度是3（3个线程），最终打印出块索引，线程索引。

并行计算

在CUDA的核函数内，每个线程通常是独立执行的，并处理对应的数据元素。在下面的案例中，每个线程负责处理矩阵中的一个特定元素，这个特定元素由线程的索引和块的索引决定。

定义8x6的矩阵：即x方向=8，y方向=6

  0  1  2  3  4  5  6  7
  8  9 10 11 12 13 14 15
 16 17 18 19 20 21 22 23
 24 25 26 27 28 29 30 31
 32 33 34 35 36 37 38 39
 40 41 42 43 44 45 46 47

定义block的维度为(4, 2)，及4x2个线程。
grid可以通过计算，矩阵是8行6列，block是(4，2)，则用6个block按照(2，3)排列即可完全覆盖数据。矩阵的行列对应block，grid的y和x。

dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);//向上取整

可以推测，8x6的矩阵将会划分成(2，3)子块，每个子块大小为(4，2)

将矩阵转置（x代表行，y代表列），如下图所示分割成（2,3）个block

在这里插入图片描述

block的维度，x方向是4，y方向是2，block的索引如下图：

在这里插入图片描述

对于第一个block，有8个线程，局部线程索引如下图所示：

在这里插入图片描述

核函数定义：下面打印局部线程id，全局线程id，以及所在block，和对应的额矩阵值

__global__ void printThreadIndex(int *A, const int nx, const int ny){
    int ix = threadIdx.x + blockIdx.x * blockDim.x;
    int iy = threadIdx.y + blockIdx.y * blockDim.y;
    unsigned int idx = iy * nx + ix;

    printf("thread_id (%d,%d) block_id (%d,%d) coordinate (%d,%d) global index %2d ival %2d\n", 
        threadIdx.x, threadIdx.y, blockIdx.x, blockIdx.y,ix, iy, idx, A[idx]);
}

threadIdx.x、threadIdx.y：线程在块内的索引。
blockIdx.x、blockIdx.y：当前块在网格中的索引。
blockDim.x、blockDim.y：每个块中线程的布局（块的大小）。
ix和iy：计算每个线程的全局坐标，表示线程在整个网格中的位置。
idx：线性索引，根据二维坐标ix、iy计算出的线程在矩阵中的全局索引。

总代码：

#include <cuda_runtime.h>
#include <stdio.h>


void printMatrix(int *C, const int nx, const int ny){
    int *ic = C;
    printf("\nMatrix: (%d.%d)\n", nx, ny);

    for (int iy = 0; iy < ny; iy++){
        for (int ix = 0; ix < nx; ix++){
            printf("%3d", ic[ix]);
        }
        ic += nx;
        printf("\n");
    }

    printf("\n");
    return;
}

__global__ void printThreadIndex(int *A, const int nx, const int ny)
{
    int ix = threadIdx.x + blockIdx.x * blockDim.x;
    int iy = threadIdx.y + blockIdx.y * blockDim.y;
    unsigned int idx = iy * nx + ix;

    printf("thread_id (%d,%d) block_id (%d,%d) coordinate (%d,%d) Matrix index %2d element %2d\n", 
            threadIdx.x, threadIdx.y, blockIdx.x, blockIdx.y,ix, iy, idx, A[idx]);
}

int main(int argc, char **argv)
{
    printf("%s Starting...\n", argv[0]);

    // get device information
    int dev = 0;
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, dev);
    printf("Using Device %d: %s\n", dev, deviceProp.name);
    cudaSetDevice(dev);

    // set matrix dimension
    int nx = 8;
    int ny = 6;
    int nxy = nx * ny;
    int nBytes = nxy * sizeof(float);

    // malloc host memory
    int *h_A;
    h_A = (int *)malloc(nBytes);

    // iniitialize host matrix with integer
    for (int i = 0; i < nxy; i++){
        h_A[i] = i;
    }
    printMatrix(h_A, nx, ny);

    // malloc device memory
    int *d_MatA;
    cudaMalloc((void **)&d_MatA, nBytes);

    // transfer data from host to device
    cudaMemcpy(d_MatA, h_A, nBytes, cudaMemcpyHostToDevice);

    // set up execution configuration
    dim3 block(4, 2);
    dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);

    // invoke the kernel
    printThreadIndex<<<grid, block>>>(d_MatA, nx, ny);
    cudaGetLastError();

    // free host and devide memory
    cudaFree(d_MatA);
    free(h_A);

    // reset device
    cudaDeviceReset();

    return (0);
}

输出：

a Starting...
Using Device 0: NVIDIA GeForce MX450

Matrix: (8.6)
  0  1  2  3  4  5  6  7
  8  9 10 11 12 13 14 15
 16 17 18 19 20 21 22 23
 24 25 26 27 28 29 30 31
 32 33 34 35 36 37 38 39
 40 41 42 43 44 45 46 47

thread_id (0,0) block_id (0,2) coordinate (0,4) Matrix index 32 element 32
thread_id (1,0) block_id (0,2) coordinate (1,4) Matrix index 33 element 33
thread_id (2,0) block_id (0,2) coordinate (2,4) Matrix index 34 element 34
thread_id (3,0) block_id (0,2) coordinate (3,4) Matrix index 35 element 35
thread_id (0,1) block_id (0,2) coordinate (0,5) Matrix index 40 element 40
thread_id (1,1) block_id (0,2) coordinate (1,5) Matrix index 41 element 41
thread_id (2,1) block_id (0,2) coordinate (2,5) Matrix index 42 element 42
thread_id (3,1) block_id (0,2) coordinate (3,5) Matrix index 43 element 43
thread_id (0,0) block_id (1,1) coordinate (4,2) Matrix index 20 element 20
thread_id (1,0) block_id (1,1) coordinate (5,2) Matrix index 21 element 21
thread_id (2,0) block_id (1,1) coordinate (6,2) Matrix index 22 element 22
thread_id (3,0) block_id (1,1) coordinate (7,2) Matrix index 23 element 23
thread_id (0,1) block_id (1,1) coordinate (4,3) Matrix index 28 element 28
thread_id (1,1) block_id (1,1) coordinate (5,3) Matrix index 29 element 29
thread_id (2,1) block_id (1,1) coordinate (6,3) Matrix index 30 element 30
thread_id (3,1) block_id (1,1) coordinate (7,3) Matrix index 31 element 31
thread_id (0,0) block_id (1,2) coordinate (4,4) Matrix index 36 element 36
thread_id (1,0) block_id (1,2) coordinate (5,4) Matrix index 37 element 37
thread_id (2,0) block_id (1,2) coordinate (6,4) Matrix index 38 element 38
thread_id (3,0) block_id (1,2) coordinate (7,4) Matrix index 39 element 39
thread_id (0,1) block_id (1,2) coordinate (4,5) Matrix index 44 element 44
thread_id (1,1) block_id (1,2) coordinate (5,5) Matrix index 45 element 45
thread_id (2,1) block_id (1,2) coordinate (6,5) Matrix index 46 element 46
thread_id (3,1) block_id (1,2) coordinate (7,5) Matrix index 47 element 47
thread_id (0,0) block_id (0,1) coordinate (0,2) Matrix index 16 element 16
thread_id (1,0) block_id (0,1) coordinate (1,2) Matrix index 17 element 17
thread_id (2,0) block_id (0,1) coordinate (2,2) Matrix index 18 element 18
thread_id (3,0) block_id (0,1) coordinate (3,2) Matrix index 19 element 19
thread_id (0,1) block_id (0,1) coordinate (0,3) Matrix index 24 element 24
thread_id (1,1) block_id (0,1) coordinate (1,3) Matrix index 25 element 25
thread_id (2,1) block_id (0,1) coordinate (2,3) Matrix index 26 element 26
thread_id (3,1) block_id (0,1) coordinate (3,3) Matrix index 27 element 27
thread_id (0,0) block_id (0,0) coordinate (0,0) Matrix index  0 element  0
thread_id (1,0) block_id (0,0) coordinate (1,0) Matrix index  1 element  1
thread_id (2,0) block_id (0,0) coordinate (2,0) Matrix index  2 element  2
thread_id (3,0) block_id (0,0) coordinate (3,0) Matrix index  3 element  3
thread_id (0,1) block_id (0,0) coordinate (0,1) Matrix index  8 element  8
thread_id (1,1) block_id (0,0) coordinate (1,1) Matrix index  9 element  9
thread_id (2,1) block_id (0,0) coordinate (2,1) Matrix index 10 element 10
thread_id (3,1) block_id (0,0) coordinate (3,1) Matrix index 11 element 11
thread_id (0,0) block_id (1,0) coordinate (4,0) Matrix index  4 element  4
thread_id (1,0) block_id (1,0) coordinate (5,0) Matrix index  5 element  5
thread_id (2,0) block_id (1,0) coordinate (6,0) Matrix index  6 element  6
thread_id (3,0) block_id (1,0) coordinate (7,0) Matrix index  7 element  7
thread_id (0,1) block_id (1,0) coordinate (4,1) Matrix index 12 element 12
thread_id (1,1) block_id (1,0) coordinate (5,1) Matrix index 13 element 13
thread_id (2,1) block_id (1,0) coordinate (6,1) Matrix index 14 element 14
thread_id (3,1) block_id (1,0) coordinate (7,1) Matrix index 15 element 15

我们取第(0,0)个block的输出，如下

thread_id (0,0) block_id (0,0) coordinate (0,0) Matrix index  0 element  0
thread_id (1,0) block_id (0,0) coordinate (1,0) Matrix index  1 element  1
thread_id (2,0) block_id (0,0) coordinate (2,0) Matrix index  2 element  2
thread_id (3,0) block_id (0,0) coordinate (3,0) Matrix index  3 element  3
thread_id (0,1) block_id (0,0) coordinate (0,1) Matrix index  8 element  8
thread_id (1,1) block_id (0,0) coordinate (1,1) Matrix index  9 element  9
thread_id (2,1) block_id (0,0) coordinate (2,1) Matrix index 10 element 10
thread_id (3,1) block_id (0,0) coordinate (3,1) Matrix index 11 element 11

第(0,0)个block的线程索引（左图）对应的矩阵值（右图）正好相对应。

在这里插入图片描述