CUDA memory troubles

Question

I have a CUDA kernel which I'm compiling to a cubin file without any special flags:

nvcc text.cu -cubin

It compiles, though with this message:

Advisory: Cannot tell what pointer points to, assuming global memory space

and a reference to a line in some temporary cpp file. I can get this to work by commenting out some seemingly arbitrary code which makes no sense to me.

The kernel is as follows:

__global__ void string_search(char** texts, int* lengths, char* symbol, int* matches, int symbolLength)
{
    int localMatches = 0;
    int blockId = blockIdx.x + blockIdx.y * gridDim.x;
    int threadId = threadIdx.x + threadIdx.y * blockDim.x;
    int blockThreads = blockDim.x * blockDim.y;

    __shared__ int localMatchCounts[32];

    bool breaking = false;
    for(int i = 0; i < (lengths[blockId] - (symbolLength - 1)); i += blockThreads)
    {
     if(texts[blockId][i] == symbol[0])
     {
      for(int j = 1; j < symbolLength; j++)
      {
       if(texts[blockId][i + j] != symbol[j])
       {
        breaking = true;
        break;
       }
      }
      if (breaking) continue;
      localMatches++;
     }
    }

    localMatchCounts[threadId] = localMatches;

    __syncthreads();

    if(threadId == 0)
    {
        int sum = 0;
        for(int i = 0; i < 32; i++)
        {
            sum += localMatchCounts[i];
        }
        matches[blockId] = sum;
    }
}

If I replace the line

localMatchCounts[threadId] = localMatches;

after the first for loop with this line

localMatchCounts[threadId] = 5;

it compiles with no notices. This can also be achieved by commenting out seemingly random parts of the loop above the line. I have also tried replacing the local memory array with a normal array to no effect. Can anyone tell me what the problem is?

The system is Vista 64bit, for what its worth.

Edit: I fixed the code so it actually works, though it still produces the compiler notice. It does not seem as though the warning is a problem, at least with regards to correctness (it might affect performance).

Answer 1

The problem seems to be associated with the char** parameter. Turning this into a char* solved the warning, so I suspect that cuda might have problems with this form of data. Perhaps cuda prefers that one uses the specific cuda 2D arrays in this case.

Answer 2

To Morten If you have passed char** to the kernel, can you tell me how you have define those memory spaces in host code and how you called the kernel??? I am trying to get a char** datatype onto the device. Im just not able to figure it out.