define LOCAL_SIZE 256

OpenCL 内存优化是提升 GPU 程序性能的关键,以下是核心优化策略和技巧：

内存层级理解

全局内存 (Global Memory) - 高延迟，所有线程可访问
常量内存 (Constant Memory) - 只读，有缓存
局部内存 (Local Memory) - 工作组内共享，低延迟
私有内存 (Private Memory) - 线程私有，寄存器级别
图像内存 (Image Memory) - 带缓存的特殊对象

关键优化策略

1 全局内存优化

// 1. 合并访问（Coalesced Access）
// 坏：分散访问
__kernel void bad_access(__global float* data) {
    int id = get_global_id(0);
    float value = data[id * 2];  // 非连续访问
}
// 好：合并访问
__kernel void good_access(__global float4* data) {
    int id = get_global_id(0);
    float4 vec = data[id];  // 连续访问，一次读4个float
}
// 2. 使用向量数据类型
__kernel void vectorized(__global float4* input, __global float4* output) {
    int id = get_global_id(0);
    float4 a = input[id];
    float4 b = a * (float4)(2.0f);
    output[id] = b;
}

2 局部内存优化

__kernel void reduction(__global float* input, __global float* output) {
    __local float local_data[LOCAL_SIZE];
    int gid = get_global_id(0);
    int lid = get_local_id(0);
    // 1. 加载到局部内存
    local_data[lid] = input[gid];
    barrier(CLK_LOCAL_MEM_FENCE);
    // 2. 在局部内存中进行归约
    for(int stride = LOCAL_SIZE/2; stride > 0; stride >>= 1) {
        if(lid < stride) {
            local_data[lid] += local_data[lid + stride];
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }
    // 3. 写回全局内存
    if(lid == 0) {
        output[get_group_id(0)] = local_data[0];
    }
}

3 常量内存优化

// 声明为常量内存
__constant float filter_weights[9] = { /* 权重值 */ };
__kernel void apply_filter(__global float* input, __global float* output) {
    // 常量内存会自动缓存，适合只读小数据
}

4 图像内存优化

// 使用图像对象代替缓冲区
__kernel void image_process(read_only image2d_t src, 
                            write_only image2d_t dst,
                            sampler_t sampler) {
    int2 coord = (int2)(get_global_id(0), get_global_id(1));
    // 自动利用纹理缓存
    float4 pixel = read_imagef(src, sampler, coord);
    pixel = pixel * 2.0f;
    write_imagef(dst, coord, pixel);
}

实用优化技巧

1 避免银行冲突（Bank Conflict）

// 坏：可能有银行冲突
__local float shared[256];
// 好：添加填充避免冲突
__local float shared_padded[256 + 1];  // +1填充

2 内存对齐

typedef struct {
    float4 position;    // 16字节对齐
    float4 normal;      // 16字节对齐
    float2 texcoord;    // 8字节对齐
} Vertex __attribute__((aligned(32)));  // 32字节对齐

3 预取数据

__kernel void prefetch_example(__global float4* data) {
    int id = get_global_id(0);
    // 预取到私有变量
    float4 a = data[id];
    float4 b = data[id + 1];
    // 计算时使用预取的数据
    float4 result = a + b;
    // ...
}

性能优化检查表

调试工具

# 使用 CodeXL、NVIDIA Visual Profiler、Intel VTune 等工具
# 检查：
# - 内存带宽利用率
# - 缓存命中率
# - 内存事务数量
# - 内存延迟隐藏情况

最佳实践

1. 小数据优先常量内存：小于几KB的只读数据
2. 频繁共享数据用局部内存：工作组内线程共享的数据
3. 合并全局内存访问：确保连续的线程访问连续的内存地址
4. 适当的工作组大小：通常是64的倍数（如64, 128, 256）
5. 使用异步拷贝：隐藏内存延迟

通过综合应用这些策略，可以显著提升 OpenCL 程序的性能，特别是在内存受限的 GPU 架构上。

标签： LOCAL_SIZE 256