Allocating array of strings in cuda

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.