mundo - Problemas de memoria CUDA

hola mundo en cuda

Question

Tengo un kernel CUDA que estoy compilando en un archivo cubin sin ningún indicador especial:

nvcc text.cu -cubin

Se compila, aunque con este mensaje:

Asesoramiento: no se puede decir a qué apunta el puntero, asumiendo el espacio de la 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 kernel 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 reemplazo la línea

localMatchCounts[threadId] = localMatches;

después del primer ciclo for 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.

Editar: arreglé el código para que realmente funcione, 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

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.

Answer 2

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 búfer continuo y dividirlo de manera que permita el acceso paralelo.
En este caso definiría una matriz 1D que contiene todas las cadenas posicionadas una tras otra y otra matriz 1D, con un tamaño de 2 * numberOfStrings que contiene el desplazamiento de cada cadena dentro de la primera matriz y su longitud:

Por 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];