Results 1 to 5 of 5

Thread: Dynamically creating 2 dimensional local memory arrays

  1. #1
    Junior Member
    Join Date
    Aug 2011
    Posts
    4

    Dynamically creating 2 dimensional local memory arrays

    In openCL you can specify the amount of local memory you want to allocate in a kernel from host code by specifing the amount of memory to allocate in a parameter for local memory with the command
    Code :
    clSetKernelArg(myKernel, 3, localHeight * localWidth * sizeof(float), NULL);
    where the kernel looks like
    Code :
    __kernel void matrixMul_gpu(
    	__global float* A, __global float* B, __global float* C,
    	__local float * As, __local float * Bs,
    	unsigned int HeightA, unsigned int WidthB, unsigned int WidthAHeightB
    	)
    This works well if I want to create 1D arrays in local memory, but what if I want to make As and Bs 2D arrays where their sizes are dynamically specified from the host code?

  2. #2
    Senior Member
    Join Date
    May 2010
    Location
    Toronto, Canada
    Posts
    845

    Re: Dynamically creating 2 dimensional local memory arrays

    You can either declare the local variable in kernel scope instead of as a kernel argument or you will have to manually index into the array as if it was 2D. This is not any different from how C99 works, is it?
    Disclaimer: Employee of Qualcomm Canada. Any opinions expressed here are personal and do not necessarily reflect the views of my employer. LinkedIn profile.

  3. #3
    Junior Member
    Join Date
    Aug 2011
    Posts
    4

    Re: Dynamically creating 2 dimensional local memory arrays

    Quote Originally Posted by david.garcia
    You can either declare the local variable in kernel scope instead of as a kernel argument or you will have to manually index into the array as if it was 2D. This is not any different from how C99 works, is it?
    Those are two work arounds that I have used.

    Declaring a local memory array within the scope of the kernel is one way to make a 2D array. But the allocation size then becomes static. I couldn't write code that dynamically sizes the array based on how much local memory is available on the devices compute units. Doing this would be useful if writing code for both nVidia's G80 and Fermi architectures. G80 only has a local memory cache of 16 kb whereas Fermi has a cache of 48 kb. If you're doing matrix dot product operations, it is ideal to have a block size as large as possible, on G80 the max would be 16x16, but on fermi it would be 32x32. If I create the array with in the scope of the kernel, I would have to write two kernels, one for a local mem size of 16x16 and the other for a local mem size of 32x32. Then I'm not even taking into account AMD cards or other architectures.

    If I try to manually index into the code, it would be more dynamic, I could specify local memory array sizes outside the kernel so I would only have to write one kernel. Unfortunetly I would have to manually index the array with something like
    Code :
    As[i * MatrixWidth + j]
    instead of
    Code :
    As[i][j]
    This would work, but it takes a few more instructions, which can add up in iterations, it can get confusing, and it makes for much less elagant code.

    I'm trying to get the best of both worlds, having dynamic code and having elagant code. Is dynamically allocating 2D arrays from the kernel call impossible in OpenCL 1.0? Or is there some way to do it?

  4. #4
    Senior Member
    Join Date
    May 2010
    Location
    Toronto, Canada
    Posts
    845

    Re: Dynamically creating 2 dimensional local memory arrays

    I couldn't write code that dynamically sizes the array based on how much local memory is available on the devices compute units.
    You can change the size dynamically. All you have to do is change the variable declaration in the source code before you call clCreateProgramWithSource().

    This would work, but it takes a few more instructions
    Why would it take more instructions? A[j][i] is identical to A[j*width+i]. As for elegance, readability, etc. you can use a macro if you want.
    Disclaimer: Employee of Qualcomm Canada. Any opinions expressed here are personal and do not necessarily reflect the views of my employer. LinkedIn profile.

  5. #5
    Senior Member
    Join Date
    May 2010
    Location
    Toronto, Canada
    Posts
    845

    Re: Dynamically creating 2 dimensional local memory arrays

    Sorry, wrong button.

    Is dynamically allocating 2D arrays from the kernel call impossible in OpenCL 1.0?
    It's not possible. See section 6.8: "Arguments to __kernel functions in a program cannot be declared as a pointer to a pointer(s)."
    Disclaimer: Employee of Qualcomm Canada. Any opinions expressed here are personal and do not necessarily reflect the views of my employer. LinkedIn profile.

Similar Threads

  1. Support of Multi-dimensional arrays
    By Sheliak in forum Suggestions for next release
    Replies: 1
    Last Post: 10-19-2011, 04:16 PM
  2. Replies: 3
    Last Post: 05-22-2011, 07:39 PM

Posting Permissions

  • You may not post new threads
  • You may not post replies
  • You may not post attachments
  • You may not edit your posts
  •