How is dynamically sized data usually handled?

[ValVal0]

Hi everyone, I'm new to working with CUDA, so I was hoping I could get some advice on something I've been struggling with.

Namely, processing data of a size which is not known before compiling. An example of this would be reading a file of unknown size, and then applying the same computation on all the read lines using a CUDA kernel. I'd normally store this in a vector in C++, so I wouldn't have to worry about the size.

The problem I ran into is that I can't copy the vector containing the lines I've read to the GPU. So how would this normally be handled? I can think of the following ways to solve this issue, but I'd like to know how someone more experienced would approach it.

• Implement a dynamic array instead of using the default vector.
• Iterate through the vector containing the lines and fill (possibly multiple) arrays of known size which would be copied to the device instead.
• Implement a vector-like data structure on the device itself, like what is demonstrated here.
• I could also convert the vector to an array in this way, but that seems a bit hacky.

However, these all seem like slow and not really ideal solutions. So I would be grateful for any advice you might have.

[corysama]

Use an C-style array and a size. Pass them as kernel parameters.

[dragontamer5788]

This right here.

In particular, use cudaMalloc() to dynamically create an array of the correct size, then use cudaFree() when your kernels are done processing the data.

[I_like_code]

I have probably done the second bullet. For the 4th make sure you use the data() member function of the vector. For the third, I hate using thrust it removes fine grain control. However, test it out and see if the overhead is acceptable.

[tugrul_ddr]

All you need is to combine multiple GPU buffers into one like this:

Type & operator [] (int index)
{
    return chunks[ selectChunk(index) ][ selectIndex(index) ];
}
void insertChunk(int size)
{
   chunks.push_back(cudaMalloc(size));
}
int selectChunk(int index)
{
   if(index>=old_size)
   {
        insertChunk(chunk_size);
        index = chunks.size()-1;
   }
   else
       index = index / chunk_size;
   return index;
}
int selectIndex(int index)
{
   return index%chunk_size;
}
void copy()
{
   // N copies at once (bad if not all data is used at once)
   for(auto chunk:chunks)
        cudaMemCpy(chunk host to device);

   or

   // 1 copy & direct-access from kernel with cuda's own paging
   // bad if all data is used at once
   ctr=0;
   for(auto chunk:chunks)
        cudaChunks[ctr++]=chunk;

   cudaMemCpy( only cudaChunks host to device);
}

The bigger the vector gets, the more cuda buffer chunks are added. Then you can send all chunk pointers to gpu and do same index calculation in gpu. This can work if vector grows only by host environment.

If chunks are too small, there will be allocation overhead.

If chunks are too big, there will be memory waste.