CUDA:具有重叠边界的共享内存分配

Question

是否有一种简单的方法(谷歌尚未提供...)从单个输入数组分配每 block 共享内存区域，以便可以重叠？

简单的例子是字符串搜索；看到我想将输入文本切 block ，让每个 block 中的每个线程搜索从文本 [thread_id] 开始的模式，但希望分配给每个 block 的数据按模式长度重叠，以便匹配跨越边界的情况还是找到了。

即分配给每个 block 上的共享内存的总内存大小是

(blocksize+patternlength)*sizeof(char)

我可能遗漏了一些简单的东西，目前正在深入阅读 CUDA 指南，但希望得到一些指导。

更新:我怀疑有些人误解了我的问题(或者我错误地解释了它)。

假设我有一个 QWERTYUIOP 数据集，我想搜索 3 个字符的匹配项，然后我将数据集(任意)切成 4 个线程 block ； QWER TYUI OPxx

这很容易完成，但如果 3 个字符匹配实际上是在寻找 IOP，则算法会失败。

在这种情况下，我想要的是每个 block 都在共享内存中:

QWERTY TYUIOP OPxxxx

即每个 block 都分配了 blocksize+patternlength-1 个字符，因此不会出现内存边界问题。

希望这能更好地解释事情。

因为@jmilloy 是持久的... :P

//VERSION 1: Simple
__global__ void gpuSearchSimple(char *T, int lenT, char *P, int lenP, int *pFound)
{
  int startIndex = blockDim.x*blockIdx.x + threadIdx.x;
    int fMatch = 1;
    for (int i=0; i < lenP; i++)
    {
      if (T[startIndex+i] != P[i]) fMatch = 0;
    }
    if (fMatch) atomicMin(pFound, startIndex);
}
//VERSION 2: Texture
__global__ void gpuSearchTexture(int lenT, int lenP, int *pFound)
{
  int startIndex = blockDim.x*blockIdx.x + threadIdx.x;
    int fMatch = 1;
    for (int i=0; i < lenP; i++)
    {
      if (tex1Dfetch(texT,startIndex+i) != tex1Dfetch(texP,i)) fMatch = 0;
    }
    if (fMatch) atomicMin(pFound, startIndex);
}
//Version 3: Shared
__global__ void gpuSearchTexSha(int lenT, int lenP, int *pFound)
{
  extern __shared__ char shaP[];
  for (int i=0;threadIdx.x+i<lenP; i+=blockDim.x){
    shaP[threadIdx.x+i]= tex1Dfetch(texP,threadIdx.x+i);
  }
  __syncthreads();

  //At this point shaP is populated with the pattern
  int startIndex = blockDim.x*blockIdx.x + threadIdx.x;
    // only continue if an earlier instance hasn't already been found
    int fMatch = 1;
    for (int i=0; i < lenP; i++)
    {
      if (tex1Dfetch(texT,startIndex+i) != shaP[i]) fMatch = 0;
    }
    if (fMatch) atomicMin(pFound, startIndex);
}

我想做的是将文本放入共享内存块中，如问题的其余部分所述，而不是将文本保留在纹理内存中以供后续版本使用。

Answer 1

我不确定这个问题是否有意义。您可以像这样在运行时动态调整共享分配内存的大小:

__global__ void kernel()
{
    extern __shared__ int buffer[];
    ....
}

kernel<<< gridsize, blocksize, buffersize >>>();

但是缓冲区的内容在内核开始时是未定义的。您将不得不在内核中设计一个方案，以从全局内存中加载您想要确保模式匹配将按您希望的方式工作的重叠。