Topic
  • 6 replies
  • Latest Post - ‏2011-04-29T13:27:41Z by SystemAdmin
SystemAdmin
SystemAdmin
131 Posts

Pinned topic is there only local memory on PS3 OpenCL

‏2011-04-28T15:10:16Z |
Hi, I have question about global, local memory and registers.
In GPU there are several types of memory, for example:
 __kernel void example(__global unsigned char *in) {
__local int localInput;
size_t i = 0;
}
On NVIDIA GPU:
leading __global indicates using of global memory
leading __local indicate using of local memory (or shared memory in CUDA model)
and size_t i = 0; is the value defined in registers.

How would it be if I this kernel execute on IBM OpnCL platform (on PS3).
Is there registers too?
Updated on 2011-04-29T13:27:41Z at 2011-04-29T13:27:41Z by SystemAdmin
  • SystemAdmin
    SystemAdmin
    131 Posts

    Re: is there only local memory on PS3 OpenCL

    ‏2011-04-28T18:14:56Z  
    On the SPE accelerator device:

    __global == system memory (accessed through a SW cache)
    __local == local storage
    __private == local storage

    The size_t i variable is allocated (if needed) in local storage memory. If it is an automatic variable, then it typically never occupies memory (local storage) and lives exclusively in registers unless dereferenced or spilled (to stack, which rarely occurs on the SPU because of its large register file).
  • SystemAdmin
    SystemAdmin
    131 Posts

    Re: is there only local memory on PS3 OpenCL

    ‏2011-04-28T21:52:41Z  
    On the SPE accelerator device:

    __global == system memory (accessed through a SW cache)
    __local == local storage
    __private == local storage

    The size_t i variable is allocated (if needed) in local storage memory. If it is an automatic variable, then it typically never occupies memory (local storage) and lives exclusively in registers unless dereferenced or spilled (to stack, which rarely occurs on the SPU because of its large register file).
    So I can rewrite this code on PS3 accelerator device like this:
     __kernel void example(__global unsigned char *in) {
    __local int localInput;
    __local size_t i = 0;
    }
    And it will be the same?

    Because when I trying to copy trough async_work_group_copy() function, from memory defined as global to memory defined like "size_t i" (according to you it's local memory), compiler gave me an error:
    "IBM_OpenCL_kernel.cl", line 50.49: 1506-324 (S) "__private struct {...}* __private" cannot be converted to "__local struct {...}*".
    Here is the code:
     __kernel void update( __global hashState *gHs) {
    int glId = get_global_id(0);
    gHs += glId;

    hashState ctx;
    async_work_group_copy((__local hashState *) &ctx, (__global hashState *) gHs, 1, NULL);
    }
    So why this error when you wrote __private == local storege == __local ?
  • SystemAdmin
    SystemAdmin
    131 Posts

    Re: is there only local memory on PS3 OpenCL

    ‏2011-04-28T22:17:18Z  
    So I can rewrite this code on PS3 accelerator device like this:
     __kernel void example(__global unsigned char *in) {
    __local int localInput;
    __local size_t i = 0;
    }
    And it will be the same?

    Because when I trying to copy trough async_work_group_copy() function, from memory defined as global to memory defined like "size_t i" (according to you it's local memory), compiler gave me an error:
    "IBM_OpenCL_kernel.cl", line 50.49: 1506-324 (S) "__private struct {...}* __private" cannot be converted to "__local struct {...}*".
    Here is the code:
     __kernel void update( __global hashState *gHs) {
    int glId = get_global_id(0);
    gHs += glId;

    hashState ctx;
    async_work_group_copy((__local hashState *) &ctx, (__global hashState *) gHs, 1, NULL);
    }
    So why this error when you wrote __private == local storege == __local ?
    Even though both __local and __private are instantiated in SPE's local storage, they still unique and separate address spaces (as specified by the OpenCL spec).

    This code will also not work in that async copy only copies arrays of OpenCL data types, not arbitrary structures.
  • SystemAdmin
    SystemAdmin
    131 Posts

    Re: is there only local memory on PS3 OpenCL

    ‏2011-04-29T06:28:03Z  
    Even though both __local and __private are instantiated in SPE's local storage, they still unique and separate address spaces (as specified by the OpenCL spec).

    This code will also not work in that async copy only copies arrays of OpenCL data types, not arbitrary structures.
    Ok, I understand, Thanks.
  • SystemAdmin
    SystemAdmin
    131 Posts

    Re: is there only local memory on PS3 OpenCL

    ‏2011-04-29T06:37:58Z  
    Even though both __local and __private are instantiated in SPE's local storage, they still unique and separate address spaces (as specified by the OpenCL spec).

    This code will also not work in that async copy only copies arrays of OpenCL data types, not arbitrary structures.
    Have one issue at the end. What if I copy data from this structure to OpenCL data types trough async_work_group_copy function. Example

     typedef struct {
    char u;
    } myStruct;
    __kernel void example( __global myStruct *in) {
    char test;
    async_work_group_copy((__global char *) &in.u, (__local char *) &test, 1, NULL);
    }
    Will this code works?
  • SystemAdmin
    SystemAdmin
    131 Posts

    Re: is there only local memory on PS3 OpenCL

    ‏2011-04-29T13:27:41Z  
    Have one issue at the end. What if I copy data from this structure to OpenCL data types trough async_work_group_copy function. Example

     typedef struct {
    char u;
    } myStruct;
    __kernel void example( __global myStruct *in) {
    char test;
    async_work_group_copy((__global char *) &in.u, (__local char *) &test, 1, NULL);
    }
    Will this code works?
    This should work as long as you wait on the event returned by async_work_group_copy. Since you are moving only a single character, a straight copy will also work.
    
    test = in->
    
    char;