Проблемы с памятью CUDA
У меня есть ядро CUDA, которое я компилирую в файл cubin без каких-либо специальных флагов:
nvcc text.cu -cubin
Он компилируется с таким сообщением:
Advisory: Cannot tell what pointer points to, assuming global memory space
и ссылка на строку в некотором временном файле cpp. Я могу заставить это работать, закомментировав какой-то, казалось бы, произвольный код, который для меня не имеет смысла.
Ядро выглядит следующим образом:
__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;
}
}
Если я заменю строку
localMatchCounts[threadId] = localMatches;
после первого цикла for с этой строкой
localMatchCounts[threadId] = 5;
он компилируется без уведомлений. Этого также можно добиться, закомментировав кажущиеся случайными части цикла над строкой. Я также попытался заменить массив локальной памяти обычным массивом, но безрезультатно. Может кто подскажет, в чем проблема?
Система Vista 64bit, чего бы она ни стоила.
Обновлено: я исправил код, чтобы он действительно работал, хотя по-прежнему выдает уведомление компилятора. Не похоже, что предупреждение является проблемой, по крайней мере, в отношении правильности (это может повлиять на производительность).
Ответы 2
Проблема, похоже, связана с параметром char **. Превращение этого в char * разрешило предупреждение, поэтому я подозреваю, что у cuda могут быть проблемы с этой формой данных. Возможно, cuda предпочитает, чтобы в этом случае использовались конкретные 2D-массивы cuda.
Массивы указателей, такие как char **, проблематичны в ядрах, поскольку ядра не имеют доступа к памяти хоста.
Лучше выделить один буфер непрерывный и разделить его таким образом, чтобы обеспечить параллельный доступ.
В этом случае я бы определил одномерный массив, который содержит все строки, расположенные одна за другой, и еще один одномерный массив размером 2 * numberOfStrings, который содержит смещение каждой строки в первом массиве и его длину:
Например - подготовка к ядру:
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]);
}
In kernel:
currentIndex = threadId + blockId * numberOfBlocks; char* currentString = buffer + metadata[2 * currentIndex]; int currentStringLength = metadata[2 * currentIndex + 1];