Problemi di memoria CUDA

Question

Ho un kernel CUDA che sto compilando in un file cubin senza flag speciali:

nvcc text.cu -cubin

Compila, sebbene con questo messaggio:

  

Avviso: impossibile dire a cosa punta il puntatore, ipotizzando spazio di memoria globale

e un riferimento a una riga in alcuni file cpp temporanei. Posso farlo funzionare commentando un codice apparentemente arbitrario che non ha senso per me.

Il kernel è il seguente:

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

Se sostituisco la linea

localMatchCounts[threadId] = localMatches;

dopo il primo ciclo per questa riga

localMatchCounts[threadId] = 5;

si compila senza preavviso. Ciò può essere ottenuto anche commentando parti apparentemente casuali del loop sopra la linea. Ho anche provato a sostituire l'array di memoria locale con un array normale senza alcun effetto. Qualcuno può dirmi qual è il problema?

Il sistema è Vista a 64 bit, per quello che vale.

Modifica: ho corretto il codice in modo che funzioni effettivamente, sebbene produca ancora l'avviso del compilatore. Non sembra che l'avvertimento sia un problema, almeno per quanto riguarda la correttezza (potrebbe influire sulle prestazioni).

Answer 1

Le matrici di puntatori come char ** sono problematiche nei kernel, poiché i kernel non hanno accesso alla memoria dell'host.
È meglio allocare un singolo buffer continuo e dividerlo in modo da consentire l'accesso parallelo.
In questo caso definirei un array 1D che contiene tutte le stringhe posizionate una dopo l'altra e un altro array 1D, dimensioni 2 * numberOfStrings che contiene l'offset di ogni stringa all'interno del primo array e la sua lunghezza:

Ad esempio - preparazione per il 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]);
}

Nel kernel:

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

Answer 2

Il problema sembra essere associato al parametro char **. Trasformarlo in un carattere * risolve l'avvertimento, quindi sospetto che cuda possa avere problemi con questa forma di dati. Forse cuda preferisce usare uno specifico array 2D cuda in questo caso.