I'm passing a struct pointer to my kernel -- how can I make that data accessible to helper functions?

Good day, folks!

I have a structure with 100+ parameters needed by my kernel. I create a buffer object for the data, do the write, and set a pointer to that 开发者_如何学Pythondata as a kernel argument. (The kernel arg is __global, but I've tried other types for it.)

So far so good! I can see the structure's elements from my kernel function just fine!

However, I want a dozen helper functions to have access to these parameters. I have tried, but I haven't found a way to do this. If I try to copy the (__global) vh into another global pointer, or a __local pointer, or a __private pointer, it fails. If I try to copy the structure data itself into a __global, or __local, or __private copy of the structure, it fails. I've tried byte-wise copies, I've tried async_work_group_copy, I've tried casts, I tried passing the whole __global pointer in to the helper functions, I've tried other types for the kernel arg itself, I've tried everything I can think of. It seems that it might not be possible to get this data to these helper functions, but it must be possible, right?

Any answer will be welcome, even if it's "that can't be done", or "you're an idiot". I've asked this question on another forum and had nobody has said even that much, though perhaps I didn't word my question properly. But ... I can't be the only person in the world with helper functions in their kernel code, right? How the heck do you get data to them that was passed in to the kernel function?

Thanks folks.... David Thanks....

---

Here is an example which demonstrates the functionality you want, hope it helps!

Structure definitions:

typedef struct agent {
    uint energy;
    uint action;
    uint type;
    uint next;
} AGENT __attribute__ ((aligned (16)));

typedef struct sim_params {
    uint size_x;
    uint size_y;
    uint size_xy;
    uint max_agents;
    uint null_agent_pointer;
    uint grass_restart;
    uint lines_per_thread;
} SIM_PARAMS;

typedef struct cell {
    uint grass;
    uint agent_pointer;
} CELL;

Helper function:

/*
 * Helper function
 */
void removeAgentFromCell(__global AGENT * agents, 
    __global CELL * matrix,
    uint cellIndex,
    uint agentIndex,
    uint previousAgentIndex,
    SIM_PARAMS sim_params) 
{

    ...
}

Main kernel:

/*
 * The kernel
 */
__kernel void step1(__global AGENT * agents, 
        __global CELL * matrix,
        __global ulong * seeds,
        const uint turn,
        const SIM_PARAMS sim_params)
{
    uint index;
    uint agentIndex;
    uint previousAgentIndex;
    ...
    // Call helper function
    removeAgentFromCell(agents, matrix, index, agentIndex, previousAgentIndex, sim_params);
    ...
}

Tried and tested and working in AMD APP SDK (on both CPU and GPU) and Nvidia CUDA Toolkit. So I guess it will work in OSX.