memoria CUDA preocupa

Question

que tienen un núcleo CUDA, que estoy compilar en un archivo cubin sin ningún tipo de marcas especiales:

nvcc text.cu -cubin 

Compila, aunque con este mensaje:

Asesor: No se puede decir a lo que apunta el puntero, asumiendo el espacio de memoria global

y una referencia a una línea en algún archivo cpp temporal. Puedo hacer que esto funcione al comentar un código aparentemente arbitrario que no tiene sentido para mí.

El núcleo es el siguiente:

__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; 
    } 
} 

Si se sustituye la línea de

localMatchCounts[threadId] = localMatches; 

después del primer bucle con esta línea

localMatchCounts[threadId] = 5; 

compila sin avisos. Esto también se puede lograr comentando partes aparentemente aleatorias del ciclo por encima de la línea. También intenté reemplazar la matriz de memoria local con una matriz normal sin ningún efecto. Alguien puede decirme cuál es el problema?

El sistema es Vista 64bit, por lo que vale la pena.

Editar: He corregido el código por lo que realmente funciona, aunque todavía produce el aviso del compilador. No parece que la advertencia sea un problema, al menos con respecto a la corrección (podría afectar el rendimiento).

Answer 1

Las matrices de punteros como char ** son problemáticas en los kernels, ya que los kernels no tienen acceso a la memoria del host.
Es mejor asignar un único buffer continuo y dividirlo de manera que permita el acceso paralelo.
En este caso Me defino una matriz 1D que contiene todas las cuerdas colocadas una tras otra y otra matriz de 1D, de tamaño 2 * numberOfStrings que contiene el desplazamiento de cada cuerda dentro de la primera matriz y el mismo de la longitud:

Para ejemplo - preparación para kernel:

 
char* buffer = st[0] + st[1] + st[2] + ....; 
int* metadata = new int[numberOfStrings * 2]; 
int lastpos = 0; 
for (int cnt = 0; cnt < 2* numberOfStrings; cnt+=2) 
{ 
    metadata[cnt] = lastpos; 
    lastpos += length(st[cnt]); 
    metadata[cnt] = length(st[cnt]); 
} 

En kernel:

 
currentIndex = threadId + blockId * numberOfBlocks; 
char* currentString = buffer + metadata[2 * currentIndex]; 
int currentStringLength = metadata[2 * currentIndex + 1]; 

Answer 2

El problema parece estar asociado con el parámetro char **. Al convertir esto en un char * resolvió la advertencia, por lo que sospecho que cuda podría tener problemas con esta forma de datos. Quizás cuda prefiere que uno use los arreglos específicos de cuda 2D en este caso.