CUDA：具有重叠边框的共享内存分配

Question

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

简单的例子是字符串搜索;看到我想骰子输入文本，让每个块中的每个线程搜索从文本[thread_id]开始的模式，但希望分配给每个块的数据重叠模式长度，以便跨越边界的匹配案例是仍然找到。

即分配给共享内存每个块的总内存大小为

(blocksize+patternlength)*sizeof(char) 

我可能失去了一些东西简单，通过CUDA引导我目前跳水，但希望一些指导。

更新：我怀疑有些人误解了我的问题（或者我误解了它）。

说我有一个数据集QWERTYUIOP，我想搜索一个3个字符的匹配，并且我将每个线程块的数据集（任意地）分成4个; QWER TYUI OPxx

这很简单，但如果3个字符匹配实际上是在寻找IOP，算法会失败。

在这种情况下，我想是每个块的共享内存有：

QWERTY TYUIOP OPxxxx

即每块被分配块大小+ 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

编号。共享内存在块中的线程之间共享，并且只能被分配给它的块访问。您不能共享两个不同块可用的内存。

据我所知，共享内存实际上驻留在多处理器上，一个线程只能从它所运行的多处理器访问共享内存。所以这是一个物理限制。 （我猜如果两个块驻留在一个mp上，则来自一个块的线程可能无法预知地访问分配给另一个块的共享内存）。

请记住，您需要将数据从全局内存显式复制到共享内存。将字符串的重叠区域复制到不重叠的共享内存是一件简单的事情。

我认为在您需要的地方获取您的数据是开发CUDA程序所需的大部分工作。我的指导意见是，你从一个解决问题的版本开始，不先使用任何共享内存。为了这个工作，你将解决你的重叠问题，共享内存的实现将变得很简单！答案被标记为正确

__global__ void gpuSearchTexSha(int lenT, int lenP, int *pFound) 
{ 
    extern __shared__ char* shared; 

    char* shaP = &shared[0]; 
    char* shaT = &shared[lenP]; 

    //copy pattern into shaP in parallel 
    if(threadIdx.x < lenP) 
     shaP[threadIdx.x] = tex1Dfetch(texP,threadIdx.x); 

    //determine texT start and length for this block 
    blockStartIndex = blockIdx.x * gridDim.x/lenT; 
    lenS = gridDim.x/lenT + lenP - 1; 

    //copy text into shaT in parallel 
    shaT[threadIdx.x] = tex1Dfetch(texT,blockStartIndex + threadIdx.x); 
    if(threadIdx.x < lenP) 
     shaP[blockDim.x + threadIdx.x] = text1Dfetch(texT,blockStartIndex + blockDim.x + threadIdx.x) 

    __syncthreads(); 

    //We have one pattern in shaP for each thread in the block 
    //We have the necessary portion of the text (with overlaps) in shaT 

    int fMatch = 1; 
    for (int i=0; i < lenP; i++) 
    { 
     if (shaT[threadIdx.x+i] != shaP[i]) fMatch = 0; 
    } 
    if (fMatch) atomicMin(pFound, blockStartIndex + threadIdx.x); 
} 

关键笔记

---

编辑2
后：

• ，我们只需要在每块共享内存模式的一个副本 - 他们都可以使用它
• 每块需要的共享内存是lenP + lenS（其中镜头是你blocksize + patternlength）
• 内核假定gridDim.x * blockDim.x =借给（同第1版）
• 我们可以复制到共享内存并行（无需for循环，如果你有足够的线程）

Answer 2

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

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

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

但缓冲区的内容是在内核的开始不确定。您必须在内核中设计一个方案，以便从全局内存中加载重叠部分，以确保您的模式匹配能够按照您的要求工作。

Answer 3

重叠的共享内存不好，每次他们想要访问共享内存中的相同地址（尽管在架构大于等于2.0的情况下，这个缓存），线程将不得不同步。

我想到的最简单的想法是复制想要重叠的文本部分。

而不是从确切chuncks全局存储器读取：

AAAA BBBB CCCC DDDD EEEE

阅读有重叠：

AAAA BBBB CCCC CCCC DDDD EEEEE