Есть ли простой способ (google еще не доставил ...) выделить области разделяемой памяти для каждого блока из одного входного массива так, чтобы их можно было перекрывать?
Простой пример - поиск строки; Увидел, что я хочу нарезать входной текст, чтобы каждый поток в каждом блоке искал шаблон, начинающийся с text [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);
}
То, что я хотел бы сделать, это поместить текст в куски общей памяти, как описано в оставшейся части вопроса, вместо того, чтобы хранить текст в памяти текстур для более поздних версий.