CODEWITHSUNDEEP
cudagpu
Programmer
Programmer

Reputation: 6753

Allocating array of strings in cuda

Let us assume that we have the following strings that we need to store in a CUDA array.

"hi there"

"this is"

"who is"

How do we declare a array on the GPU to do this. I tried using C++ strings but it does not work.

Upvotes: 1

Views: 5556

Answers (2)

kroneml
kroneml

Reputation: 687

You have to use C-style character strings char *str. Searching for "CUDA string" on google would have given you this CUDA "Hello World" example as first hit: http://computer-graphics.se/hello-world-for-cuda.html There you can see how to use char*-strings in CUDA. Be aware that standard C-functions like strcpy or strcmp are not available in CUDA!

If you want an array of strings, you just have to use char** (as in C/C++). As for strcmp and similar functions, it highly depends on what you want to do. CUDA is not really well suited for string operations, maybe it would help if you would provide a little more detail about what you want to do.

Upvotes: -1

talonmies
talonmies

Reputation: 72349

Probably the best way to do this is to use structure that is similar to common compressed sparse matrix formats. Store the character data packed into a single piece of linear memory, then use a separate integer array to store the starting indices, and perhaps a third array to store the string lengths. The storage overhead of the latter might be more efficient that storing a string termination byte for every entry in the data and trying to parse for the terminator inside the GPU code.

So you might have something like this:

struct gpuStringArray {
    unsigned int * pos; 
    unsigned int * length;  // could be a smaller type if strings are short
    char4 * data; // 32 bit data type will improve memory throughput, could be 8 bit
}

Note I used a char4 type for the string data; the vector type will give better memory throughput, but it will mean strings need to be aligned/suitably padded to 4 byte boundaries. That may or may not be a problem depending on what a typical real string looks like in your application. Also, the type of the (optional) length parameter should probably be chosen to reflect the maximum admissible string length. If you have a lot of very short strings, it might be worth using an 8 or 16 bit unsigned type for the lengths to save memory.

A really simplistic code to compare strings stored this way in the style of strcmp might look something like this:

__device__ __host__
int cmp4(const char4 & c1, const char4 & c2)
{
    int result;

    result = c1.x - c2.x; if (result !=0) return result; 
    result = c1.y - c2.y; if (result !=0) return result; 
    result = c1.z - c2.z; if (result !=0) return result; 
    result = c1.w - c2.w; if (result !=0) return result; 

    return 0;
}

__device__ __host__
int strncmp4(const char4 * s1, const char4 * s2, const unsigned int nwords)
{
    for(unsigned int i=0; i<nwords; i++) {
        int result = cmp4(s1[i], s2[i]);
        if (result != 0) return result;
    }

    return 0;
}

__global__
void tkernel(const struct gpuStringArray a, const gpuStringArray b, int * result)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;

    char4 * s1 = a.data + a.pos[idx];
    char4 * s2 = b.data + b.pos[idx];
    unsigned int slen = min(a.length[idx], b.length[idx]);

    result[idx] = strncmp4(s1, s2, slen);
}

[disclaimer: never compiled, never tested, no warranty real or implied, use at your own risk]

There are some corner cases and assumptions in this which might catch you out depending on exactly what the real strings in your code look like, but I will leave those as an exercise to the reader to resolve. You should be able to adapt and expand this into whatever it is you are trying to do.

Upvotes: 2

Related Questions