CUDA内存问题

Question

我有一个CUDA内核，我正在将其编译为cubin文件，没有任何特殊的标志：

nvcc text.cu -cubin

编译成功，但是出现以下信息:

警告: 无法确定指针指向，假定为全局内存空间

并且引用了某个临时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 64位，如果有用的话。

编辑：我修复了代码，使其实际工作，虽然它仍然会产生编译器提示。至少在正确性方面（可能会影响性能），警告似乎不是问题。

- Morten Christiansen

Answer 1

指针数组（如char**）在内核中存在问题，因为内核无法访问主机的内存。
最好是分配一个单一的连续缓冲区，并以一种使并行访问成为可能的方式进行划分。
在这种情况下，我会定义一个1D数组，其中包含所有字符串依次放置，另一个1D数组，大小为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]);
}

在内核中：

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

- Danny

Answer 2

问题似乎与char **参数有关。将其转换为char *解决了警告，因此我怀疑cuda可能会在处理这种形式的数据时出现问题。也许在这种情况下，cuda更喜欢使用特定的cuda 2D数组。

- Morten Christiansen