Page 1 of 2 12 LastLast
Results 1 to 10 of 16

Thread: passing an array of char* to kernel function

  1. #1
    Junior Member
    Join Date
    Oct 2010
    Posts
    26

    passing an array of char* to kernel function

    hello,
    I'm an OpenCL beginner and not a C++ professional, I've read the oclVectorAdd example in the OpenCL samples... I found that the arguments of the kernel function were float * (float pointers representing 3 arrays of floats)... I've got an idea of a GPU program in which I want to create a simple kernel function but this time I want to deal with an array of strings... since there are no strings the cl program I'm trying to send 2 arrays of char* as an argument to the kernel function... Just for testing the idea, I wrote this code:

    Code :
    __kernel void normalizer(__global char * documentStrings[], __global char * documentStringsOut[])
    { int i = get_global_id(0);
    int strLen=sizeof(documentStrings[i])/sizeof(char);
    for(int j=0;j<strLen;j++){
        char ch = documentStrings[i][j]; documentStrings[i][j]=ch++;}
    documentStringsOut[i]=documentStrings[i];
    }

    But when I try to call the clBuildProgram I get this error: "invalid address space for array argument to __kernel function"

    can anybody help me?!

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

    Re: passing an array of char* to kernel function

    The reason you are seeing an error is because you are declaring "documentStrings" and "documentStringsOut" as arrays of pointers to char (i.e. "char *foo[]"). Arrays of pointers are not allowed as kernel arguments.

    There are ways to accomplish the sort of thing you are trying to do but they are rather complicated for a beginner.

    If I were you I would try to implement the code first in plain C (with the restriction above about arrays of pointers not allowed) and only once you have the code working then think about porting it to OpenCL.
    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
    Member
    Join Date
    Oct 2010
    Location
    Vancouver, Canada
    Posts
    65

    Re: passing an array of char* to kernel function

    I agree with David's suggestion to try this in plain C first, that will simplify your development work substantially. I will give you an additional hint as to how I could approach this though -- since you can't use pointers, the usual fallback is to use indices. Consider having a large buffer/array of chars, and instead of pointing at each string, have an index into the array to the start of your string.

    Once you have the indices and the array in your kernel, you can switch to pointers if you prefer, but working with indices can have advantages.

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

    Re: passing an array of char* to kernel function

    Thanks, Andrew. I hesitated on whether to give that advice or not

    Once you have the indices and the array in your kernel, you can switch to pointers if you prefer, but working with indices can have advantages.
    Keep in mind that in OpenCL you can't work with pointers in the device's address space. I would stick to indices for good.
    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
    Junior Member
    Join Date
    Oct 2010
    Posts
    26

    Re: passing an array of char* to kernel function

    Thanks a lot Mr. David & Mr. Andrew for your help! I got your idea and I'll try implementing it but I tried another idea too (it might be less practical than yours, most probably I'll go for yours, but I just want to understand)
    The other idea I thought of was to wrap the char* in a structure then passing a pointer to this structure to the kernel function... the code was as follows:
    Code :
    typedef struct
    { char* c;
    }string_t;
     
     __kernel void normalizer(__global string_t* documentStrings, __global string_t* documentStringsOut){ int i = get_global_id(0); int strLen=sizeof(documentStrings[i].c)/sizeof(char);	
    for(int j=0;j<strLen;j++){	char ch = documentStrings[i].c[j]; documentStrings[i].c[j]=ch++;}
     documentStringsOut[i].c=documentStrings[i].c;
    }

    It ran without errors but it behaved in a strange way! the returned array (documentsStringsOut) had only one element, the first string only and the rest were "Bad Ptr"!
    I wish to understand why? I implemented it on an array of 5 strings and in the clEnqueueNDRangeKernel function I set the global_work_size to 5 and the work_dim to 1...

    I'm sorry for my long message but I really want to learn and understand... thanks very much for your help

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

    Re: passing an array of char* to kernel function

    Code :
    typedef struct
    { char* c;
    }string_t;

    That's an illegal struct declaration in OpenCL and the compiler should have produced a compilation error.

    You also want to give a look at this code:
    Code :
    int strLen=sizeof(documentStrings[i].c)/sizeof(char);

    Sizeof documentStrings[i].c doesn't mean "the length of this string". It means "the size of a pointer to char", which is typically 4 or 8 bytes depending on your platform.

    If you want to compute the length of a null-terminated string you need to use the function strlen() from the C standard library, which unfortunately doesn't exist in OpenCL.

    Spend some time practicing your C first and it will save you a lot of frustration when you try OpenCL
    Disclaimer: Employee of Qualcomm Canada. Any opinions expressed here are personal and do not necessarily reflect the views of my employer. LinkedIn profile.

  7. #7
    Member
    Join Date
    Mar 2010
    Location
    Raleigh, NC
    Posts
    55

    Re: passing an array of char* to kernel function

    Quote Originally Posted by david.garcia
    Spend some time practicing your C first and it will save you a lot of frustration when you try OpenCL
    This statement is so incredibly true that it no longer is funny to me. I'm a grad student in aerospace engineering (not a native programmer), and started with Fortran, and now am working a lot with C++ and Fortran. There's a lot of "hand-holding" that occurs in these languages that don't exist in C, and so when I write my OpenCL programs, I end up having to tweak a good many little C++ things that slip into my C kernels.

    I am getting better...david.garcia is a wonderful asset on these forums!

  8. #8
    Member
    Join Date
    Oct 2010
    Location
    Vancouver, Canada
    Posts
    65

    Re: passing an array of char* to kernel function

    An important thing to realize about pointers in OpenCL is that they differ between devices. If you are running on the CPU device that happens to be the same underlying hardware as the host, then you may happen to find that the pointers are identical... but this is a very bad assumption to rely on, and it will break badly on any other devices (and isn't guaranteed by the spec to ever work, even if your host is the exact same hardware as your device).

    OpenCL devices may have independent address spaces from the host (and each other). There is no assurance that the physical bytes in a buffer when accessed from the host are the same as the physical bytes in the same buffer when accessed from any other device, or that the physical bytes will be the same ones holding the buffer over time. Even if they do happen to be the same physical bytes, they could be mapped to different addresses. When you pass a cl_mem object to a kernel, the system is responsible for converting that object handle into the device's native pointers. This may include physically copying the data between memory spaces (for example, from host memory to a GPU's VRAM). Since the bits that make up pointers look just like the bits that make up any other data, the OpenCL runtime has no idea what bits are pointers that need to be remapped... except when they are kernel arguments. Thus you cannot put pointers into your buffers and expect them to work from anywhere else (even a different work-group or kernel on the same device).

    Even beyond having different addresses, you aren't even assured that pointers are the same number of bits. If your host is running a 64-bit OS then its pointers are going to be 64-bit if your application is 64-bit. Most GPU's, however, are not 64-bit devices and thus will use 32-bit pointers... or smaller. The different address spaces of a device (global, local, constant, private) can also vary in size and pointers could potentially be only 16-bits or even smaller for really small memory pools.

    It is also important to understand the memory model that OpenCL provides -- it is a "relaxed consistency memory model" and this means that there are very carefully defined rules about when memory needs to be made consistent between devices, kernels, work-groups, work-items ... and these rules are defined to allow as low a degree of synchronization and coupling as possible, thus maximizing opportunity for concurrency. This means you need to be careful about having multiple pieces of code touching the same buffer at the same time... even if it seems to work on your particular machine when you write the code, any number of factors can cause your assumptions to break (changing hardware, changing OS, changing drivers, and even just changes to the way your own application works). This affects things like when writes to a buffer from one piece of code become visible to reads from the buffer in a different work-item, kernel, device, etc. Or what the result is when multiple pieces of code write to the same location "simultaneously" (and I put that in quotes because "simultaneous" is a very hard thing to define and harder to be assured of).


    So this all boils down to: do not make assumptions about pointers and do not pass them between devices. The CL memory objects were put into the spec to abstract these details from us. Respect and understand that abstraction.

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

    Re: passing an array of char* to kernel function

    Amen to everything Andrew just said. He put it very eloquently.
    Disclaimer: Employee of Qualcomm Canada. Any opinions expressed here are personal and do not necessarily reflect the views of my employer. LinkedIn profile.

  10. #10
    Junior Member
    Join Date
    Oct 2010
    Posts
    26

    Re: passing an array of char* to kernel function

    I don't know how to thank you Andrew for the detailed explanation... this made things much more clear.... by the way I tried implementing your idea of the character buffer and working on indices instead of pointers... it's now working...
    Also I guess I'll stick to your advice David, of practicing C first... and I'd appreciate any advice of how to improve my self in C and OpenCl...
    Thanks again.

Page 1 of 2 12 LastLast

Similar Threads

  1. Replies: 3
    Last Post: 01-30-2013, 03:34 AM
  2. Replies: 10
    Last Post: 11-08-2011, 06:28 AM

Posting Permissions

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