Create local array dynamic inside OpenCL kernel

Question

I have a OpenCL kernel that needs to process a array as multiple arrays where each sub-array sum is saved in a local cache array.

For example, imagine the fowling array:

[[1, 2, 3, 4], [10, 30, 1, 23]]

Each work-group gets a array (in the exemple we have 2 work-groups);

Each work-item process two array indexes (for example multiply the value index the local_id), where the work-item result is saved in a work-group shared array.

__kernel void test(__global int **values, __global int *result, const int array_size){
    __local int cache[array_size];

    // initialise
    if (get_local_id(0) == 0){
        for (int i = 0; i < array_size; i++)
            cache[i] = 0;
    }

    barrier (CLK_LOCAL_MEM_FENCE);

    if(get_global_id(0) < 4){
        for (int i = 0; i<2; i++)
            cache[get_local_id(0)] += values[get_group_id(0)][i] * 
                                                         get_local_id(0);
    }

    barrier (CLK_LOCAL_MEM_FENCE);

    if(get_local_id(0) == 0){
        for (int i = 0; i<array_size; i++)
            result[get_group_id(0)] += cache[i];
    }
}

The problem is that I can not define the cache array size by using a kernel parameter, but i need to in order to have a dynamic kernel.

How can I create it dynamically? like malloc function in c...

Or the only solution available is to send a temp array to my kernel function?

stackoverflow.com/questions/2541929/…

Ciro Santilli OurBigBook.com
– Ciro Santilli OurBigBook.com

2016-03-31 16:25:31 +00:00
Commented Mar 31, 2016 at 16:25 — Ciro Santilli OurBigBook.com
– Ciro Santilli OurBigBook.com, Commented Mar 31, 2016 at 16:25

Dmitry Shkuropatsky · Accepted Answer · 2013-07-11 00:22:56Z

18

This can be achieved by adding __local array as a kernel parameter:

__kernel void test(__global int **values, __global int *result, 
    const int array_size, __local int * cache)

and providing desired size of the kernel parameter:

clSetKernelArg(kernel, 3, array_size*sizeof(int), NULL);

The local memory will be allocated upon the kernel invocation. Note, that extra checks may be necessary to ensure that required local memory size does not exceed the device limit.

edited Jul 11, 2013 at 0:22

answered Jul 10, 2013 at 16:40

Dmitry Shkuropatsky

3,9752 gold badges23 silver badges13 bronze badges

Sign up to request clarification or add additional context in comments.

Collectives™ on Stack Overflow

Create local array dynamic inside OpenCL kernel

1 Answer 1

Comments

Your Answer

Linked

Hot Network Questions

Collectives™ on Stack Overflow

1 Answer 1

Comments

Your Answer

Sign up or log in

Post as a guest

Linked

Related