OpenCL 結構體資料傳遞與核函數定義完整指南

核心原則：宿主端與裝置端定義必須完全一致

OpenCL 允許在 Kernel 中使用自定義結構體，但要求宿主端 (Host) 與裝置端 (Device) 的結構體記憶體佈局完全相同。這包括：欄位順序、資料類型、對齊方式、填充位元組。任何不一致都會導致資料解讀錯誤。

宿主端定義範例

// 宿主端結構體定義
typedef struct StudentNode {
    int age;        // 4 bytes
    float height;   // 4 bytes
    // 總計 8 bytes，無填充 (假設 4-byte 對齊)
} Student;

// 記憶體分配與初始化
Student* stu_input = (Student*)malloc(sizeof(Student));
stu_input->age = 25;
stu_input->height = 1.80f;

Student* stu_output = (Student*)malloc(sizeof(Student));

建立 Buffer 並傳遞給 Kernel

// 輸入 Buffer：複製宿主端資料
cl_mem input_buffer = clCreateBuffer(context,
    CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
    sizeof(Student), stu_input, &err);

// 輸出 Buffer：僅分配裝置端記憶體
cl_mem output_buffer = clCreateBuffer(context,
    CL_MEM_WRITE_ONLY,
    sizeof(Student), NULL, &err);

// 設定 Kernel 參數
cl_kernel kernel = clCreateKernel(program, "structTest", &err);
clSetKernelArg(kernel, 0, sizeof(cl_mem), &input_buffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &output_buffer);

裝置端 Kernel 定義

// 必須與宿主端完全一致的結構體定義
typedef struct StudentNode {
    int age;
    float height;
} Student;

__kernel void structTest(__global const Student* input,
                         __global Student* output) {
    // 直接存取結構體欄位
    output->age = input->age + 1;
    output->height = input->height * 1.1f;
}

進階：結構體陣列與動態大小

// 宿主端：結構體陣列
typedef struct Particle {
    float4 position;
    float4 velocity;
    float mass;
    int type;
} Particle;

const int NUM_PARTICLES = 10000;
Particle* particles = (Particle*)malloc(NUM_PARTICLES * sizeof(Particle));
// 初始化粒子資料...

// Buffer 建立
cl_mem particle_buffer = clCreateBuffer(context,
    CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
    NUM_PARTICLES * sizeof(Particle), particles, &err);

// Kernel 端
typedef struct Particle {
    float4 position;
    float4 velocity;
    float mass;
    int type;
} Particle;

__kernel void nbody_sim(__global Particle* particles, int num_particles, float dt) {
    int gid = get_global_id(0);
    if (gid >= num_particles) return;
    
    // 直接操作結構體陣列元素
    particles[gid].position += particles[gid].velocity * dt;
}

對齊與填充陷阱 (關鍵！)

最佳實踐建議

1. 共用標頭檔：將結構體定義放在 .h 檔，宿主端 #include，Kernel 端同樣 #include (需確保編譯器支援)
2. 使用標準類型：優先使用 cl_int、cl_float、cl_float4 等 OpenCL 內建類型，保證跨平台一致性
3. 避免指標成員：結構體內絕不能包含指標，裝置端無法存取宿主端位址
4. 驗證大小：執行時檢查 sizeof(Student) 在宿主端與裝置端是否一致
5. 對齊屬性：必要時使用 __attribute__((aligned(16))) 強制對齊

完整可編譯範例

// common.h (宿主端與 Kernel 共用)
#ifndef COMMON_H
#define COMMON_H

typedef struct Vec3 {
    float x, y, z;
} Vec3;

typedef struct Particle {
    Vec3 pos;
    Vec3 vel;
    float mass;
    int id;
} Particle;

#endif

// host.c
#include "common.h"
#include <CL/cl.h>
// ... 建立 context, queue, program ...
cl_mem buf = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
                            n * sizeof(Particle), particles, &err);

// kernel.cl
#include "common.h"
__kernel void update(__global Particle* p, int n, float dt) {
    int i = get_global_id(0);
    if (i < n) p[i].pos.x += p[i].vel.x * dt;
    // ...
}