有没有一种简单的方法（google还没有提供……）从单个输入数组中分配每个块的共享内存区域，这样就可以有重叠？
简单的例子是字符串搜索;我想把输入文本切成小块，让每个块中的每个线程搜索从text[thread_id]开始的模式，但是希望分配给每个块的数据按照模式长度重叠，这样仍然可以找到跨越边界的匹配情况。
即分配给每个块上的共享内存的总内存大小为
(blocksize+patternlength)*sizeof(char)我可能错过了一些简单的东西，目前正在通过CUDA指南潜水，但会感谢一些指导。
更新：我怀疑有些人误解了我的问题（或者我没有解释清楚）。
假设我有一个数据集QWERTYUIOP，我想搜索一个3个字符的匹配，我将数据集（任意）切成4个线程块; QWER TYUI OPxx
这很容易实现，但如果3个字符匹配实际上是在查找IOP，则算法会失败。
在这种情况下，我想要的是每个块都在共享内存中：
QWERTY TYUIOP OPxxxx即，每个块被分配X1 M4 N1 X字符，因此不会发生存储器边界问题。
希望这能更好地解释事情。
因为@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);
}我想做的是将文本放入共享内存块中，如问题的其余部分所述，而不是将文本保存在纹理内存中以供以后的版本使用。
3条答案
按热度按时间voj3qocg1#
我不确定这个问题有多大意义。您可以在运行时动态调整共享分配内存的大小，如下所示：
但是缓冲区的内容在内核的开始处是未定义的。您必须在内核中设计一个方案，以便从全局内存中加载您想要的重叠部分，以确保模式匹配将按照您想要的方式工作。
uxhixvfz2#
不可以。共享内存是在一个块中的线程之间共享的，并且只能被分配给它的块访问。不能有可用于两个不同块的共享内存。
据我所知，共享内存实际上驻留在多处理器上，线程只能从它所运行的多处理器访问共享内存。这是一个物理限制。（我猜如果两个块驻留在一个mp上，来自一个块的线程可能无法访问分配给另一个块的共享内存）。
请记住，您需要显式地将数据从全局内存复制到共享内存。将字符串的重叠区域复制到不重叠的共享内存是一件简单的事情。
我认为在需要的地方获取数据是开发CUDA程序所需的大部分工作。我的指导是，您可以从一个解决问题的版本开始，而不首先使用任何共享内存。为了使其工作，您将解决重叠问题，并且共享内存实现将很容易！
编辑2
关键注解：
lenP + lenS（其中透镜为您的blocksize + patternlength）mnemlml83#
重叠共享内存并不好，线程每次想要访问共享内存中的同一地址时都必须同步（尽管在架构>= 2.0中，这已经得到了缓解）。
我想到的最简单的想法是复制你想要重叠的文本部分。
而不是从全局内存中读取精确的块：
AAAA BBBB CCCC DDDD EEEE
重叠读取：
AAAA BBBB CCCCCCCCDDDD EEEEE