CUDA详解

NiNg_1_234

910人浏览 · 2025-01-11 17:14:07

NiNg_1_234 · 2025-01-11 17:14:07 发布

文章目录

CUDA详解

CUDA详解

在这里插入图片描述

一、引言

CUDA（Compute Unified Device Architecture）是由NVIDIA推出的并行计算平台和编程模型，它允许开发者利用NVIDIA GPU的强大计算能力来加速计算密集型任务。CUDA通过扩展C/C++语言，提供了一种高效的方式来编写并行程序，广泛应用于科学计算、深度学习和高性能计算等领域。

二、CUDA编程模型

1、线程层次结构

CUDA的编程模型基于线程层次结构，主要包括线程（Thread）、线程块（Thread Block）和网格（Grid）三个层次：

线程：最基本的执行单位，每个线程执行相同的代码，但处理不同的数据。
线程块：线程的集合，同一块内的线程可以共享数据和同步操作。
网格：线程块的集合，网格内的线程块独立执行，互不影响。

2、内存层次结构

CUDA的内存层次结构包括寄存器（Register）、共享内存（Shared Memory）、全局内存（Global Memory）等：

寄存器：每个线程有自己的寄存器，访问速度最快。
共享内存：同一线程块内的线程共享，访问速度较快。
全局内存：所有线程共享，访问速度较慢。

三、CUDA编程步骤

1、定义核函数

核函数（Kernel Function）是在GPU上并行执行的函数，使用__global__关键字定义。例如：

__global__ void addKernel(int *c, const int *a, const int *b)
{
    int i = threadIdx.x;
    c[i] = a[i] + b[i];
}

2、内存分配与数据传输

在主机（CPU）和设备（GPU）之间进行内存分配和数据传输：

int main(void)
{
    int a[256], b[256], c[256];
    int *dev_a, *dev_b, *dev_c;

    // 分配设备内存
    cudaMalloc((void**)&dev_a, 256 * sizeof(int));
    cudaMalloc((void**)&dev_b, 256 * sizeof(int));
    cudaMalloc((void**)&dev_c, 256 * sizeof(int));

    // 初始化主机数据
    for (int i = 0; i < 256; i++) {
        a[i] = -i;
        b[i] = i * i;
    }

    // 将数据从主机复制到设备
    cudaMemcpy(dev_a, a, 256 * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, 256 * sizeof(int), cudaMemcpyHostToDevice);
}

3、启动核函数

使用<<< >>>语法启动核函数，指定网格和线程块的大小：

addKernel<<<1, 256>>>(dev_c, dev_a, dev_b);

4、同步与内存回收

等待核函数执行完毕，并回收设备内存：

cudaDeviceSynchronize();

// 将结果从设备复制回主机
cudaMemcpy(c, dev_c, 256 * sizeof(int), cudaMemcpyDeviceToHost);

// 释放设备内存
cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_c);

四、使用示例

1、矩阵乘法

矩阵乘法是一个典型的并行计算任务，可以利用CUDA高效实现。以下是一个简单的矩阵乘法示例：

#define BLOCK_SIZE 16

__global__ void matrixMul(float *C, float *A, float *B, int width)
{
    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;
    int Row = by * BLOCK_SIZE + ty;
    int Col = bx * BLOCK_SIZE + tx;
    float Cvalue = 0;

    for (int i = 0; i < width; i += BLOCK_SIZE)
    {
        __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
        __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];

        As[ty][tx] = A[Row * width + i + tx];
        Bs[ty][tx] = B[(i + ty) * width + Col];
        __syncthreads();

        for (int k = 0; k < BLOCK_SIZE; ++k)
            Cvalue += As[ty][k] * Bs[k][tx];
        __syncthreads();
    }

    C[Row * width + Col] = Cvalue;
}

int main()
{
    int width = 256;
    float *A, *B, *C;
    float *dev_A, *dev_B, *dev_C;

    // 分配主机内存
    A = (float*)malloc(width * width * sizeof(float));
    B = (float*)malloc(width * width * sizeof(float));
    C = (float*)malloc(width * width * sizeof(float));

    // 初始化矩阵
    for (int i = 0; i < width * width; i++) {
        A[i] = 1.0f;
        B[i] = 2.0f;
    }

    // 分配设备内存
    cudaMalloc((void**)&dev_A, width * width * sizeof(float));
    cudaMalloc((void**)&dev_B, width * width * sizeof(float));
    cudaMalloc((void**)&dev_C, width * width * sizeof(float));

    // 将数据从主机复制到设备
    cudaMemcpy(dev_A, A, width * width * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_B, B, width * width * sizeof(float), cudaMemcpyHostToDevice);

    // 启动核函数
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
    dim3 dimGrid((width + BLOCK_SIZE - 1) / BLOCK_SIZE, (width + BLOCK_SIZE - 1) / BLOCK_SIZE);
    matrixMul<<<dimGrid, dimBlock>>>(dev_C, dev_A, dev_B, width);

    // 同步
    cudaDeviceSynchronize();

    // 将结果从设备复制回主机
    cudaMemcpy(C, dev_C, width * width * sizeof(float), cudaMemcpyDeviceToHost);

    // 释放内存
    free(A);
    free(B);
    free(C);
    cudaFree(dev_A);
    cudaFree(dev_B);
    cudaFree(dev_C);

    return 0;
}