1
// 寄存器变量：默认情况下，函数内的局部变量都存储在寄存器中
2
__global__ void kernel() {
3
    float x = 1.0f;        // x存储在寄存器中
4
    int i = threadIdx.x;   // i存储在寄存器中
5
    // 寄存器是线程私有的，每个线程都有自己的副本
6
}

1
// 共享内存：用 __shared__ 关键字声明
2
__global__ void kernel_with_shared(int *data) {
3
    __shared__ int shared_data[256];  // 每个block有一份共享内存
4
    int tid = threadIdx.x;
5

6
    // 从全局内存加载到共享内存
7
    shared_data[tid] = data[tid];
8
    __syncthreads();  // 等待所有线程完成加载
9

10
    // 所有线程可以访问共享内存中的数据
11
    int sum = shared_data[0] + shared_data[255];
12
}

1
// 全局内存：用 cudaMalloc 分配
2
__global__ void kernel(int *d_data) {
3
    // d_data指向全局内存
4
    int val = d_data[threadIdx.x];  // 数百周期的延迟！
5
}

1
Grid（线程网格）
2
  └── Block（线程块）
3
        └── Thread（线程）

1
// 查询设备限制
2
cudaDeviceProp prop;
3
cudaGetDeviceProperties(&prop, 0);
4
printf("Max threads per block: %d\n", prop.maxThreadsPerBlock);
5
printf("Max grid size: %d x %d x %d\n",
6
       prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
7
printf("Shared memory per block: %zu KB\n", prop.sharedMemPerBlock / 1024);

1
kernel_name<<<grid_dim, block_dim, shared_mem_size, stream>>>(args);

1
// 1D配置：处理N个元素
2
int N = 1024 * 1024;
3
int threads_per_block = 256;
4
int blocks = (N + threads_per_block - 1) / threads_per_block;  // 向上取整
5
kernel<<<blocks, threads_per_block>>>(d_data, N);
6

7
// 2D配置：处理图像
8
dim3 block(16, 16);   // 16x16 = 256线程
9
dim3 grid((width + 15) / 16, (height + 15) / 16);
10
kernel_2d<<<grid, block>>>(d_image, width, height);

1
float *d_data;
2
size_t bytes = N * sizeof(float);
3

4
// 在GPU上分配内存
5
cudaError_t err = cudaMalloc(&d_data, bytes);
6
if (err != cudaSuccess) {
7
    printf("cudaMalloc failed: %s\n", cudaGetErrorString(err));
8
}
9

10
// 使用完毕后释放
11
cudaFree(d_data);

1
float *h_data = (float*)malloc(bytes);  // 主机内存
2
float *d_data;
3
cudaMalloc(&d_data, bytes);
4

5
// 主机 → 设备（同步）
6
cudaMemcpy(d_data, h_data, bytes, cudaMemcpyHostToDevice);
7

8
// 启动kernel
9
kernel<<<grid, block>>>(d_data, N);
10

11
// 设备 → 主机（同步）
12
cudaMemcpy(h_result, d_data, bytes, cudaMemcpyDeviceToHost);

1
// 1. 分配页锁定内存（主机端）
2
float *h_data;
3
cudaMallocHost(&h_data, bytes);  // 页锁定，可被DMA直接访问
4

5
// 2. 分配设备内存
6
float *d_data;
7
cudaMalloc(&d_data, bytes);
8

9
// 3. 创建Stream
10
cudaStream_t stream;
11
cudaStreamCreate(&stream);
12

13
// 4. 异步传输 + 异步Kernel启动（都在同一个stream中）
14
cudaMemcpyAsync(d_data, h_data, bytes, cudaMemcpyHostToDevice, stream);
15
kernel<<<grid, block, 0, stream>>>(d_data, N);
16
cudaMemcpyAsync(h_result, d_data, bytes, cudaMemcpyDeviceToHost, stream);
17

18
// 5. 等待stream完成
19
cudaStreamSynchronize(stream);
20

21
// 6. 清理
22
cudaStreamDestroy(stream);
23
cudaFreeHost(h_data);
24
cudaFree(d_data);

1
时间轴（单Stream）：
2
H2Dchunk1 | Kernelchunk1 | D2Hchunk1 | H2Dchunk2 | Kernelchunk2 | D2Hchunk2
3

4
时间轴（双Stream重叠）：
5
Stream1: H2Dchunk1 | Kernelchunk1 | D2Hchunk1
6
Stream2:            H2Dchunk2 | Kernelchunk2 | D2Hchunk2
7
                      ↑ 重叠区域：Stream1计算时，Stream2在传输

1
const int STREAM_COUNT = 4;
2
cudaStream_t streams[STREAM_COUNT];
3
for (int i = 0; i < STREAM_COUNT; i++) {
4
    cudaStreamCreate(&streams[i]);
5
}
6

7
int chunk_size = N / STREAM_COUNT;
8
for (int i = 0; i < STREAM_COUNT; i++) {
9
    int offset = i * chunk_size;
10
    cudaMemcpyAsync(&d_data[offset], &h_data[offset],
11
                    chunk_size * sizeof(float),
12
                    cudaMemcpyHostToDevice, streams[i]);
13
    kernel<<<grid, block, 0, streams[i]>>>(&d_data[offset], chunk_size);
14
    cudaMemcpyAsync(&h_result[offset], &d_data[offset],
15
                    chunk_size * sizeof(float),
16
                    cudaMemcpyDeviceToHost, streams[i]);
17
}
18

19
// 等待所有stream完成
20
for (int i = 0; i < STREAM_COUNT; i++) {
21
    cudaStreamSynchronize(streams[i]);
22
    cudaStreamDestroy(streams[i]);
23
}