Блог

Есть ли простой способ (Google не предоставил …) выделить области общей памяти для каждого блока из одного входного массива таким образом, чтобы могло быть перекрытие?

Простым примером является поиск строки; Увидел, что я хочу разделить входной текст на части, заставить каждый поток в каждом блоке искать шаблон, начинающийся с text[thread_id] , но хочу, чтобы данные, назначенные каждому блоку, перекрывались по длине шаблона, чтобы совпадающие случаи, которые попадают через границу, все еще были найдены.

Т.е. общий объем памяти, выделенный для общей памяти в каждом блоке, равен

 (blocksize patternlength)*sizeof(char)
  

Вероятно, я упускаю что-то простое и в настоящее время просматриваю руководство по CUDA, но был бы признателен за некоторые рекомендации.

ОБНОВЛЕНИЕ: я подозреваю, что некоторые люди неправильно поняли мой вопрос (или я пропустил его объяснение).

Допустим, у меня есть набор данных QWERTYUIOP , и я хочу выполнить поиск соответствия из 3 символов, и я разбиваю набор данных (произвольно) на 4 для каждого блока потока; QWER TYUI OPxx

Это достаточно просто выполнить, но алгоритм завершается неудачей, если действительно ищется совпадение из 3 символов IOP .

В этом случае я хочу, чтобы каждый блок находился в общей памяти:

 QWERTY TYUIOP OPxxxx
  

т.е. каждому блоку присваиваются blocksize patternlength-1 символы, поэтому проблем с границей памяти не возникает.

Надеюсь, это объясняет ситуацию лучше.

Поскольку @jmilloy является постоянным … 😛

 //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);
}
  

Что я хотел бы сделать, так это поместить текст в блоки общей памяти, как описано в остальной части вопроса, вместо сохранения текста в памяти текстур для более поздних версий.

Я не уверен, что этот вопрос имеет такой уж большой смысл. Вы можете динамически изменять размер совместно выделяемой памяти во время выполнения следующим образом:

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

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

но содержимое буфера не определено в начале ядра. Вам нужно будет разработать схему в ядре для загрузки из глобальной памяти с перекрытием, которое вы хотите, чтобы гарантировать, что ваше сопоставление с шаблоном будет работать так, как вы хотите.

Нет. Общая память распределяется между потоками в блоке и доступна только тому блоку, которому она назначена. У вас не может быть общей памяти, доступной для двух разных блоков.

Насколько я знаю, общая память фактически находится в мультипроцессорах, и поток может получить доступ к общей памяти только из мультипроцессора, на котором он запущен. Итак, это физическое ограничение. (Я предполагаю, что если два блока находятся на одном mp, поток из одного блока может иметь возможность непредсказуемо обращаться к общей памяти, которая была выделена другому блоку).

Помните, что вам необходимо явно скопировать данные из глобальной памяти в общую память. Скопировать перекрывающиеся области строки в неперекрывающуюся общую память несложно.

Я думаю, что получение ваших данных там, где они вам нужны, составляет большую часть работы, необходимой при разработке программ CUDA. Я советую вам начать с версии, которая решает проблему без использования какой-либо общей памяти в первую очередь. Для того, чтобы это сработало, вы решите свою проблему перекрытия, и реализация общей памяти станет простой!

правка 2
после того, как ответ был отмечен как правильный

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

    char* shaP = amp;shared[0];
    char* shaT = amp;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);
}
  

основные примечания:

• нам нужна только одна копия шаблона в общей памяти на блок — они все могут ее использовать
• для каждого блока требуется общая память lenP lenS (где lenS — ваш blocksize patternlength )
• ядро предполагает, что gridDim.x * blockDim.x = предоставлено (то же, что и в версии 1)
• мы можем копировать в общую память параллельно (циклы for не нужны, если у вас достаточно потоков)

Перекрывающаяся общая память — это плохо, потоку придется синхронизировать каждый раз, когда они захотят получить доступ к одному и тому же адресу в общей памяти (хотя в архитектуре > = 2.0 это было смягчено).

Самая простая идея, которая приходит мне в голову, — это продублировать ту часть текста, которую вы хотите перекрыть.

Вместо чтения из глобальной памяти точными порциями:

ААААА BBBBB CCCC DDDD ИИИИ

Чтение с перекрывающимися границами:

ААААА BBBBB CCCC CCCC DDDD EEEEE