CUDA 入门¶

约 856 个字 101 行代码预计阅读时间 4 分钟

这并不是一个严格意义上的 CUDA 入门，只是希望帮助大家更好的去理解 CUDA 代码的并行性以及其运行在 GPU 上这个事实

与运行在 CPU 上的 C 代码区别¶

GPU 有大量核心，天生具有并行性，因此作为 GPU 的编程语言，CUDA 的编程更类似于并行编程。

CUDA 函数主要特点如下:

我们需要指定这个函数使用的线程数目 func<<<gridDim, blockDim>>>(...)，这样指定后，GPU 便会将这这些线程分布到其众多核心上并行执行，CUDA 函数所干的事情就是安排好每个线程做什么。一般来说，若干线程组成一个 block，若干 block 又组成一个 grid，总线程数目就是 blockDim * gridDim（注：blockDim 和 gridDim 可以是三维的，为了解释方便我们就认为其是一个数了）。
由于有很多个线程同时执行这个函数，线程与线程之间的执行顺序是完全随机的，我们不能假设线程 0，线程 1 之间可以同时执行完某一部分代码，也不能假定他们有什么特定的先后关系。这就需要我们在 CUDA 函数中添加同步点，来保证每个线程都可以看到正确的结果。
其只能使用 CUDA 指针与特定类型的函数，这个也非常好理解，GPU 能直接访问的只有它自己的显存，CPU 上的内存是不好直接用的。同时如果存在 CUDA 内部的函数调用，它也只能调用带有 device 属性的 CUDA 函数以及带有 global 属性的 CUDA 函数，不能调用 C 函数。
代码中在调用 CUDA 函数后，CPU 不会等待 CUDA 函数执行完毕，而是会直接继续执行后面的代码，如果我们没有进行同步，直接调用函数结果的话结果就会出错。

除了这些，CUDA 函数的编写方式基本与 C 函数相差无几。　　

小试牛刀¶

下面的代码使用 GPU 实现了对数组求和的功能，但是其中有些小 bug 需要大家来找一下。程序正确运行之后可以对比两个计时点得到的计时，体会正常的 C 函数与 CUDA 函数的不同。

编译方式：

nvcc reduce.cu -o reduce

如果代码看不懂可以问问开源的模型

// reduce.cu
#include <cuda_runtime.h>
#include <iostream>
#include <time.h>



__global__ void reductionKernel(float* input, float* output, int size) {
    extern __shared__ float sharedData[];

    int tid = threadIdx.x;
    int index = blockIdx.x * blockDim.x + threadIdx.x;

    // Load data into shared memory
    sharedData[tid] = (index < size) ? input[index] : 0;
    // __syncthreads();

    // Perform reduction in shared memory
    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tid < stride) {
            sharedData[tid] += sharedData[tid + stride];
        }
        // __syncthreads();
    }

    // Write result for this block to global memory
    if (tid == 0) {
        output[blockIdx.x] = sharedData[0];
    }
}

void checkCudaError(cudaError_t err, const char* msg) {
    if (err != cudaSuccess) {
        std::cerr << "CUDA Error: " << msg << " - " << cudaGetErrorString(err) << std::endl;
        exit(EXIT_FAILURE);
    }
}

void reduction_gpu(float* input, float* output, int size) {
    int blockSize = 256; // Number of threads per block
    int gridSize = (size + blockSize - 1) / blockSize;

    float* d_input, * d_output;
    checkCudaError(cudaMalloc(&d_input, size * sizeof(float)), "Failed to allocate device memory for input");
    checkCudaError(cudaMalloc(&d_output, gridSize * sizeof(float)), "Failed to allocate device memory for output");

    checkCudaError(cudaMemcpy(d_input, input, size * sizeof(float), cudaMemcpyHostToDevice), "Failed to copy data to device");

    // Launch the reduction kernel using cudaLaunchKernel
    clock_t start = clock();
    void* kernelArgs[] = { &input, &output, &size };
    // void* kernelArgs[] = { &d_input, &d_output, &size };
    checkCudaError(cudaLaunchKernel((void*)reductionKernel, 
                                    dim3(gridSize), dim3(blockSize), 
                                    kernelArgs, blockSize * sizeof(float), 0), 
                                    "Failed to launch reduction kernel");

    clock_t end = clock();
    double duration_kernel = (double)(end - start) / (CLOCKS_PER_SEC);
    printf("reduction_kernel: %.9lf\n", duration_kernel);

    // Check for any errors during kernel execution
    checkCudaError(cudaGetLastError(), "Kernel execution failed");
    // Wait for the kernel to complete 
    cudaDeviceSynchronize();

    end = clock();
    duration_kernel = (double)(end - start) / (CLOCKS_PER_SEC);
    printf("reduction_kernel after sync: %.9lf\n", duration_kernel);  
    // Copy the partial results back to the host
    checkCudaError(cudaMemcpy(output, d_output, gridSize * sizeof(float), cudaMemcpyDeviceToHost), "Failed to copy data back to host");
    // Perform the final reduction on the host
    float finalSum = 0.0f;
    for (int i = 0; i < gridSize; ++i) {
        finalSum += output[i];
    }

    std::cout << "Total sum: " << finalSum << std::endl;

    cudaFree(d_input);
    cudaFree(d_output);
}

int main() {
    const int size = 1024 * 1024 * 16;
    float* input = new float[size];
    float* output = new float[(size + 255) / 256];

    // Initialize input data
    for (int i = 0; i < size; ++i) {
        input[i] = static_cast<float>(2);
    }

    reduction_gpu(input, output, size);

    delete[] input;
    delete[] output;

    return 0;
}