How to create variable sized __local memory in pyopencl?

Question

in my C OpenCL code I use clSetKernelArg to create 'variable size' __local memory for use in my kernels, which is not available in OpenCL per se. See my example:

clSetKernelArg(clKernel, ArgCounter++, sizeof(cl_mem), (void *)&d_B);
...
clSetKernelArg(clKernel, ArgCounter++, sizeof(float)*block_size*block_size, NULL);
...
kernel="    
    matrixMul(__global float* C,
                   ...
              __local float* A_temp,
                  ...
             )"
{...

My question is now, how to do the same in pyopencl?

I looked through the examples that come with pyopencl, but the only thing I could find was an approach using templates, which seems as to me as I understood it like an overkill. See example.

kernel = """
        __kernel void matrixMul(__global float* C,...){
                     ... 
            __local float A_temp[ %(mem_size) ];
                     ...
        }

What do you recommend?

Chad Brewbaker · Accepted Answer · 2011-05-06 22:11:44Z

6

It is similar to C. You pass it a fixed size array as a local. Here is an example from Enja's radix sort. Notice the last argument is a local memory array.

 def naive_scan(self, num):
    nhist = num/2/self.cta_size*16
    global_size = (nhist,)
    local_size = (nhist,)
    extra_space = nhist / 16 #NUM_BANKS defined as 16 in RadixSort.cpp
    shared_mem_size = self.uintsz * (nhist + extra_space)
    scan_args = (   self.mCountersSum,
                    self.mCounters,
                    np.uint32(nhist),
                    cl.LocalMemory(2*shared_mem_size)
                )
    self.radix_prg.scanNaive(self.queue, global_size, local_size, *(scan_args)).wait()

answered May 6, 2011 at 22:11

Chad Brewbaker

2,5752 gold badges20 silver badges26 bronze badges

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

Comments

Rick-Rainer Ludwig · Accepted Answer · 2011-05-06 17:57:55Z

3

I am no familiar with Python and its OpenCL implementation, but a local memory can also be created within the kernel with a fixed size (similar what you did):

__kernel void matrixMul(...) {

    __local float A_templ[1024];

}

Instead of 1024 a defined preprocessor symbol can be used and can be set during compilation to change the size:

#define SIZE 1024    
__kernel void matrixMul(...) {

        __local float A_templ[SIZE];

    }

SIZE can be defined within the same soure, as compiler parameter for cLBuildProgram or as an additional source for clCreateProgramWithSource.

EDIT: Found something with Google ;-): http://www.google.com/url?sa=t&source=web&cd=4&ved=0CC8QFjAD&url=http%3A%2F%2Flinksceem.eu%2Fjoomla%2Ffiles%2FPRACE_Winter_School%2FLinkSCEMM_pyOpenCL.pdf&rct=j&q=Pyopencl%20__local%20memory&ei=BTbETbWhOsvBswadp62ODw&usg=AFQjCNG6rXEEkDpE1304pmQDu3GFdRA0BQ&sig2=vHOGOqwA1HHUl10c6HO8WQ&cad=rja

edited May 6, 2011 at 17:57

answered May 6, 2011 at 17:51

Rick-Rainer Ludwig

2,3992 gold badges27 silver badges42 bronze badges

1 Comment

Framester Over a year ago

Hi RR, I accepted Chads answer, because it is directly what I was looking for. Even though, one can find the answer in the pdf you linked to on page 21.

Collectives™ on Stack Overflow

How to create variable sized __local memory in pyopencl?

2 Answers 2

Comments

1 Comment

Your Answer

Linked

Hot Network Questions

Collectives™ on Stack Overflow

2 Answers 2

Comments

1 Comment

Your Answer

Sign up or log in

Post as a guest

Linked

Related