// OpenCL C 内核函数
__kernel void my_kernel(
    __global float *input,      // 全局内存
    __global float *output,     // 全局内存
    __constant float *weights,  // 常量内存
    __local float *scratch,     // 局部内存
    const int count             // 标量参数（通过值传递）
) {
    // ...
}

// C++ 主机端设置参数
cl_kernel kernel = clCreateKernel(program, "my_kernel", &err);

// 内存对象参数：传递 cl_mem 的地址
cl_mem buffer = clCreateBuffer(...);
clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer);

// 标量参数：传递值的地址
int count = 1024;
clSetKernelArg(kernel, 4, sizeof(int), &count);

// 局部内存：只分配大小，不传数据
clSetKernelArg(kernel, 3, 256 * sizeof(float), NULL);

// 设置完参数后，内核可以多次执行（参数不变的情况下）

// ✅ 合并访问：线程 i 访问元素 i
__kernel void good_access(__global float *data) {
    int gid = get_global_id(0);
    float val = data[gid];  // 连续访问
}

// ❌ 步长访问：线程 i 访问元素 i*stride
__kernel void bad_access(__global float *data, int stride) {
    int gid = get_global_id(0);
    float val = data[gid * stride];  // 非连续，缓存不友好
}

// ❌ 冲突：所有线程访问同一银行
float val = data[0];

// ✅ 无冲突：不同线程访问不同银行
float val = data[get_local_id(0)];

// 使用向量类型提高内存吞吐
__kernel void vectorized(__global float4 *input, __global float4 *output) {
    int gid = get_global_id(0);
    float4 val = input[gid];  // 一次加载 16 字节
    val = val * 2.0f;
    output[gid] = val;
}

__kernel void local_example(
    __global float *input,
    __global float *output,
    __local float *scratch  // 局部内存（主机端传 NULL）
) {
    int lid = get_local_id(0);
    int gid = get_global_id(0);

    // 1. 从全局内存加载到局部内存
    scratch[lid] = input[gid];
    barrier(CLK_LOCAL_MEM_FENCE);  // 同步工作组

    // 2. 在局部内存上操作（快速）
    float sum = 0.0f;
    for (int i = 0; i < get_local_size(0); i++) {
        sum += scratch[i];
    }

    // 3. 写回全局内存
    if (lid == 0) {
        output[get_group_id(0)] = sum;
    }
}

// 工作组内并行求和
__kernel void parallel_sum(
    __global float *input,
    __global float *output,
    __local float *scratch
) {
    int lid = get_local_id(0);
    int size = get_local_size(0);

    scratch[lid] = input[get_global_id(0)];
    barrier(CLK_LOCAL_MEM_FENCE);

    for (int stride = size / 2; stride > 0; stride >>= 1) {
        if (lid < stride) {
            scratch[lid] += scratch[lid + stride];
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }

    if (lid == 0) {
        output[get_group_id(0)] = scratch[0];
    }
}

cl_ulong local_mem_size;
clGetDeviceInfo(device, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(local_mem_size), &local_mem_size, NULL);
printf("Local memory: %lu KB\n", local_mem_size / 1024);  // 通常 32-64 KB

// OpenCL C 中声明
__constant float weights[256];  // 只读，64 KB 限制

__kernel void use_weights(__global float *data, __global float *out) {
    int gid = get_global_id(0);
    float sum = 0.0f;
    for (int i = 0; i < 256; i++) {
        sum += data[gid * 256 + i] * weights[i];
    }
    out[gid] = sum;
}

// 主机端：创建常量缓冲区
cl_mem weights_buf = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                                     sizeof(weights), weights_data, &err);

// 绑定到内核的 __constant 参数
clSetKernelArg(kernel, 1, sizeof(cl_mem), &weights_buf);

__kernel void private_example(__global float *data) {
    int gid = get_global_id(0);

    // 以下变量默认在私有内存（寄存器）
    float temp = data[gid];
    int count = 100;
    float results[4];  // 小数组可能在寄存器中

    // 如果寄存器溢出，编译器会自动溢出到"私有内存"（实际上是局部/全局内存）
    float large_array[1024];  // ⚠️ 可能导致寄存器压力，影响占用率
}

// 通常选择 64、128、256、512、1024
size_t local_size = 256;

// 必须是 2 的幂（某些设备要求）
// 必须 ≤ CL_DEVICE_MAX_WORK_GROUP_SIZE
// 必须能整除 global_size

// ❌ 差：工作项发散
if (gid % 2 == 0) {
    // 路径 A
} else {
    // 路径 B → GPU 执行两路
}

// ✅ 好：让工作项走同一路径
// 数据预处理，将偶数/奇数分到不同工作组

// ❌ 差：运行时除法
int row = gid / width;

// ✅ 好：如果 width 是 2 的幂，使用位运算
int row = gid >> log2_width;
int col = gid & (width - 1);

// matrix_mul.cl - 利用局部内存的矩阵乘法
#define TILE_SIZE 16

__kernel void matrix_mul(
    __global const float *A,
    __global const float *B,
    __global float *C,
    const int M,    // A 的行数
    const int N,    // B 的列数
    const int K     // A 的列数 = B 的行数
) {
    int row = get_global_id(0);
    int col = get_global_id(1);

    int lid_row = get_local_id(0);
    int lid_col = get_local_id(1);

    __local float tileA[TILE_SIZE][TILE_SIZE];
    __local float tileB[TILE_SIZE][TILE_SIZE];

    float sum = 0.0f;

    // 分块计算
    for (int t = 0; t < K; t += TILE_SIZE) {
        // 协作加载到局部内存
        tileA[lid_row][lid_col] = A[row * K + t + lid_col];
        tileB[lid_row][lid_col] = B[(t + lid_row) * N + col];
        barrier(CLK_LOCAL_MEM_FENCE);

        // 计算当前块的部分积
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += tileA[lid_row][k] * tileB[k][lid_col];
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }

    if (row < M && col < N) {
        C[row * N + col] = sum;
    }
}