CUDA 编程入门：从 Hello GPU 到线程网格模型

int globalIdx = blockIdx.x * blockDim.x + threadIdx.x;

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

// 错误检查宏——生产环境必备
#define CUDA_CHECK(call)                                                    \
    do {                                                                    \
        cudaError_t err = call;                                             \
        if (err != cudaSuccess) {                                           \
            fprintf(stderr, "CUDA error at %s:%d: %s\n",                     \
                    __FILE__, __LINE__, cudaGetErrorString(err));           \
            exit(EXIT_FAILURE);                                             \
        }                                                                   \
    } while (0)

// 核函数：每个线程计算 C[i] = A[i] + B[i]
__global__ void vectorAdd(const float *A, const float *B, float *C, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        C[i] = A[i] + B[i];
    }
}

int main() {
    const int N = 1 << 20;          // 1,048,576 个元素
    size_t bytes = N * sizeof(float);

    // 1. 分配主机内存并初始化
    float *h_A = (float *)malloc(bytes);
    float *h_B = (float *)malloc(bytes);
    float *h_C = (float *)malloc(bytes);
    for (int i = 0; i < N; i++) {
        h_A[i] = 1.0f;
        h_B[i] = 2.0f;
    }

    // 2. 分配设备内存
    float *d_A, *d_B, *d_C;
    CUDA_CHECK(cudaMalloc((void **)&d_A, bytes));
    CUDA_CHECK(cudaMalloc((void **)&d_B, bytes));
    CUDA_CHECK(cudaMalloc((void **)&d_C, bytes));

    // 3. 将数据从主机拷贝到设备（H2D）
    CUDA_CHECK(cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice));

    // 4. 启动核函数
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    CUDA_CHECK(cudaGetLastError());  // 检查核函数启动错误
    CUDA_CHECK(cudaDeviceSynchronize());

    // 5. 将结果从设备拷贝回主机（D2H）
    CUDA_CHECK(cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost));

    // 6. 验证结果
    for (int i = 0; i < 5; i++) {
        printf("C[%d] = %.1f\n", i, h_C[i]);
    }

    // 7. 释放资源
    CUDA_CHECK(cudaFree(d_A));
    CUDA_CHECK(cudaFree(d_B));
    CUDA_CHECK(cudaFree(d_C));
    free(h_A); free(h_B); free(h_C);

    return 0;
}

nvcc -o vector_add vector_add.cu -O3
./vector_add

vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

#define TILE_WIDTH 16

__global__ void matrixMulShared(const float *A, const float *B, float *C,
                                int width) {
    __shared__ float s_A[TILE_WIDTH][TILE_WIDTH];
    __shared__ float s_B[TILE_WIDTH][TILE_WIDTH];

    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;

    int row = by * TILE_WIDTH + ty;
    int col = bx * TILE_WIDTH + tx;

    float Cvalue = 0.0f;

    for (int m = 0; m < (width + TILE_WIDTH - 1) / TILE_WIDTH; m++) {
        // 协作加载 A 和 B 的子块到共享内存
        int aCol = m * TILE_WIDTH + tx;
        int bRow = m * TILE_WIDTH + ty;
        s_A[ty][tx] = (row < width && aCol < width) ? A[row * width + aCol] : 0.0f;
        s_B[ty][tx] = (bRow < width && col < width) ? B[bRow * width + col] : 0.0f;

        __syncthreads();  // 等待全部加载完成

        // 计算子块乘积
        for (int k = 0; k < TILE_WIDTH; k++) {
            Cvalue += s_A[ty][k] * s_B[k][tx];
        }

        __syncthreads();  // 等待计算完成再加载下一轮
    }

    if (row < width && col < width) {
        C[row * width + col] = Cvalue;
    }
}

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

__global__ void scaleKernel(float *data, float scale, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        data[i] *= scale;
    }
}

int main() {
    const int N = 1 << 20;
    size_t bytes = N * sizeof(float);
    int numStreams = 4;

    // 分配页锁定内存（允许异步传输）
    float *h_data;
    cudaMallocHost(&h_data, bytes);

    float *d_data;
    cudaMalloc(&d_data, bytes);

    // 创建多个流
    cudaStream_t streams[numStreams];
    for (int i = 0; i < numStreams; i++) {
        cudaStreamCreate(&streams[i]);
    }

    int chunkSize = N / numStreams;
    int threadsPerBlock = 256;
    int blocksPerGrid = (chunkSize + threadsPerBlock - 1) / threadsPerBlock;

    // 分块异步执行：拷贝 → 计算 → 拷贝回
    for (int i = 0; i < numStreams; i++) {
        int offset = i * chunkSize;
        size_t chunkBytes = chunkSize * sizeof(float);

        // 异步 H2D 拷贝
        cudaMemcpyAsync(d_data + offset, h_data + offset,
                        chunkBytes, cudaMemcpyHostToDevice, streams[i]);

        // 异步启动核函数
        scaleKernel<<<blocksPerGrid, threadsPerBlock, 0, streams[i]>>>
            (d_data + offset, 2.0f, chunkSize);

        // 异步 D2H 拷贝
        cudaMemcpyAsync(h_data + offset, d_data + offset,
                        chunkBytes, cudaMemcpyDeviceToHost, streams[i]);
    }

    // 等待所有流完成
    for (int i = 0; i < numStreams; i++) {
        cudaStreamSynchronize(streams[i]);
    }

    // 清理
    for (int i = 0; i < numStreams; i++) {
        cudaStreamDestroy(streams[i]);
    }
    cudaFree(d_data);
    cudaFreeHost(h_data);

    return 0;
}

# 查看本机 GPU 的 Compute Capability
nvidia-smi
/usr/local/cuda/extras/demo_suite/deviceQuery