GPU
├── Streaming Multiprocessors (SMs)
│   ├── CUDA Cores
│   ├── Special Function Units (SFUs)
│   ├── Load/Store Units
│   └── Shared Memory
├── Memory Controllers
├── Global Memory (DRAM)
└── L2 Cache

__global__ void kernelName(float* input, float* output, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        output[tid] = input[tid] * 2.0f;
    }
}

Grid (网格)
├── Block 0 (线程块)
│   ├── Thread 0-1023
│   └── ...
├── Block 1
│   └── ...
└── Block N

// 启动256个线程块，每个线程块512个线程
kernelName<<<256, 512>>>(d_input, d_output, n);

// 线程索引
threadIdx.x, threadIdx.y, threadIdx.z
// 块索引
blockIdx.x, blockIdx.y, blockIdx.z
// 块维度
blockDim.x, blockDim.y, blockDim.z
// 网格维度
gridDim.x, gridDim.y, gridDim.z

// 计算全局线程ID
int globalId = blockIdx.x * blockDim.x + threadIdx.x;

__global__ void kernel() {
    int temp = 10; // 寄存器变量
    // ...
}

__global__ void kernel() {
    __shared__ float sdata[256]; // 共享内存声明
    // 线程块内协作处理
    sdata[threadIdx.x] = global_data[blockIdx.x * blockDim.x + threadIdx.x];
    __syncthreads(); // 同步线程块内所有线程
}

// 在主机端分配全局内存
float *d_data;
cudaMalloc(&d_data, size * sizeof(float));

// 访问全局内存
__global__ void kernel(float* d_data) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float value = d_data[idx]; // 从全局内存读取
    // ...
}

__constant__ float const_data[1024]; // 常量内存声明

// 在主机端设置常量内存
cudaMemcpyToSymbol(const_data, h_data, size);

// 良好的合并访问
__global__ void goodCoalesced(float* input, float* output, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    output[tid] = input[tid] * 2.0f;
}

// 不好的非合并访问
__global__ void badCoalesced(float* input, float* output, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    output[tid] = input[tid * 16]; // 跳跃访问
}

__global__ void matrixMultiply(float* A, float* B, float* C, int N) {
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];

    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;

    for (int k = 0; k < N / BLOCK_SIZE; k++) {
        As[ty][tx] = A[row * N + k * BLOCK_SIZE + tx];
        Bs[ty][tx] = B[(k * BLOCK_SIZE + ty) * N + col];
        __syncthreads();

        // 计算乘积
        float sum = 0;
        for (int i = 0; i < BLOCK_SIZE; i++) {
            sum += As[ty][i] * Bs[i][tx];
        }
        C[row * N + col] = sum;
        __syncthreads();
    }
}

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

// 核函数（在GPU上执行）
__global__ void add(float* a, float* b, float* c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    int n = 1024;
    size_t size = n * sizeof(float);

    // 主机端数据
    float *h_a = (float*)malloc(size);
    float *h_b = (float*)malloc(size);
    float *h_c = (float*)malloc(size);

    // 设备端数据
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, size);
    cudaMalloc(&d_b, size);
    cudaMalloc(&d_c, size);

    // 数据传输：主机到设备
    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);

    // 启动核函数
    int blockSize = 256;
    int gridSize = (n + blockSize - 1) / blockSize;
    add<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);

    // 数据传输：设备到主机
    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);

    // 清理资源
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);

    return 0;
}

#define CUDA_CHECK(call) \
    do { \
        cudaError_t error = call; \
        if (error != cudaSuccess) { \
            fprintf(stderr, "CUDA error at %s:%d - %s\n", \
                    __FILE__, __LINE__, cudaGetErrorString(error)); \
            exit(EXIT_FAILURE); \
        } \
    } while (0)

// 使用示例
CUDA_CHECK(cudaMalloc(&d_ptr, size));
CUDA_CHECK(cudaMemcpy(d_ptr, h_ptr, size, cudaMemcpyHostToDevice));
kernel<<<grid, block>>>(d_ptr);
CUDA_CHECK(cudaGetLastError()); // 检查kernel启动错误
CUDA_CHECK(cudaDeviceSynchronize()); // 等待kernel完成

// 创建多个流
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// 异步内存拷贝和kernel执行
cudaMemcpyAsync(d_a1, h_a1, size, cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(d_a2, h_a2, size, cudaMemcpyHostToDevice, stream2);

kernel<<<grid1, block1, 0, stream1>>>(d_a1, d_b1, d_c1);
kernel<<<grid2, block2, 0, stream2>>>(d_a2, d_b2, d_c2);

cudaMemcpyAsync(h_c1, d_c1, size, cudaMemcpyDeviceToHost, stream1);
cudaMemcpyAsync(h_c2, d_c2, size, cudaMemcpyDeviceToHost, stream2);

// 同步所有流
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

// 清理资源
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);

// 查询设备属性
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);

// 计算最佳block size
int minGridSize, blockSize;
cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize,
                                  myKernel, 0, 0);

// 启动优化后的kernel
myKernel<<<minGridSize, blockSize>>>(/* parameters */);

// 展开前
__global__ void unroll1(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    for (int i = 0; i < 4; i++) {
        data[idx * 4 + i] *= 2.0f;
    }
}

// 展开后
__global__ void unroll2(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int base = idx * 4;
    data[base] *= 2.0f;
    data[base + 1] *= 2.0f;
    data[base + 2] *= 2.0f;
    data[base + 3] *= 2.0f;
}

// 使用min避免分支
__global__ void optimized(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;

    for (int i = idx; i < n; i += stride) {
        data[i] = sqrt(data[i]);
    }
}

// 使用float4进行矢量化访问
__global__ void vectorizedLoad(float4* input, float4* output, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float4 data = input[idx];
        data.x *= 2.0f;
        data.y *= 2.0f;
        data.z *= 2.0f;
        data.w *= 2.0f;
        output[idx] = data;
    }
}

// 使用__ldg读取只读数据
__global__ void readOnlyCache(const float* __restrict__ input,
                             float* output, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float val = __ldg(&input[idx]); // 使用只读缓存
        output[idx] = val * 2.0f;
    }
}

// 子kernel
__global__ void childKernel(float* data, int offset, int size) {
    int idx = threadIdx.x;
    if (idx < size) {
        data[offset + idx] *= 2.0f;
    }
}

// 父kernel
__global__ void parentKernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    // 动态启动子kernel
    childKernel<<<1, 256>>>(data, idx * 256, 256);
    cudaDeviceSynchronize(); // 同步子kernel
}

// 统一内存分配
float *unified_data;
cudaMallocManaged(&unified_data, size * sizeof(float));

// 直接访问，无需显式拷贝
kernel<<<grid, block>>>(unified_data, size);

// CPU也可以直接访问
for (int i = 0; i < size; i++) {
    unified_data[i] = i * 2.0f;
}

cudaFree(unified_data);

cudaGraph_t graph;
cudaGraphExec_t graphExec;

// 创建图
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
kernel<<<grid1, block1, 0, stream>>>(/* params */);
cudaMemcpyAsync(d_ptr, h_ptr, size, cudaMemcpyHostToDevice, stream);
kernel<<<grid2, block2, 0, stream>>>(/* params */);
cudaStreamEndCapture(stream, &graph);

// 实例化图
cudaGraphInstantiate(&graphExec, graph, NULL, NULL, 0);

// 重复执行图
for (int i = 0; i < iterations; i++) {
    cudaGraphLaunch(graphExec, stream);
    cudaStreamSynchronize(stream);
}

# 分析kernel性能
nv-nsight-cu ./my_application

# 指定特定kernel
nv-nsight-cu --kernel=my_kernel ./my_application

# 生成系统级跟踪
nsys profile ./my_application

# 查看结果
nsys-ui profile.nsys-rep

cudaProfilerStart();
kernel<<<grid, block>>>(/* parameters */);
cudaProfilerStop();

// 检查编译时常量
#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 600
    // 使用新特性
#else
    // 兼容旧架构
#endif

// 模板化不同数据类型
template<typename T>
__global__ void genericKernel(T* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        data[idx] = process(data[idx]);
    }
}

// 编译时启用调试
nvcc -g -G my_code.cu

// 使用printf调试（谨慎使用，影响性能）
__global__ void debugKernel(float* data) {
    printf("Thread %d: data = %f\n", threadIdx.x, data[threadIdx.x]);
}

// 使用assert进行运行时检查
__global__ void kernelWithAssert(int* data) {
    int idx = threadIdx.x;
    assert(idx < 1024);
    data[idx] = idx;
}