CUDA: как читать 4 (или 16) символов за одну транзакцию на поток, используя текстуры и char4 (или int4)?

У меня есть большой массив символов в глобальной памяти устройства, к которому потоки обращаются объединенным образом. Я где-то читал, что могу ускорить доступ к памяти, прочитав 4 или 16 символов в одной транзакции памяти на поток. Я считаю, что мне придется использовать текстуры и структуры char4 или int4. Однако я не могу найти никакой документации или примеров по этому поводу. Может ли кто-нибудь здесь привести простой пример или указать, где я могу узнать больше об этом?

В моем коде я определяю массив символов как

char *database = NULL;
cudaMalloc( (void**) &database, SIZE * sizeof(char) );

Каким будет определение, если я хочу использовать текстуры и char4 (или int4)?

Большое спасибо.


cuda textures texture-packing
person Ross    schedule 02.08.2012    source источник
comment
Если вы передаете deviceptr в базу данных и гарантируете, что он выровнен по 128 битам, вы можете просто прочитать его, используя int4*. Это приведет к 128-битному чтению на поток через L1.   -  person Greg Smith    schedule 03.08.2012
comment
Хорошо, я попробую это. Если SIZE кратен 16, то будет ли он выровнен по 128-битам? Это может показаться глупым вопросом, но я хочу убедиться, что правильно понял. Не могли бы вы уточнить просто прочитать? Спасибо.   -  person Ross    schedule 03.08.2012

Ответы (1)


arrow_upward
1
arrow_downward

Наконец-то я нашел ответ на свой вопрос. Определение с char4 будет

char4 *database = NULL;
cudaMalloc( (void**) &database, SIZE * sizeof(char4)/4 );

Текстуры для этого не нужны. Ядро ускоряется в три раза с char4, но уменьшается до двух, если я выполняю развертывание цикла. Для полноты картины мое ядро

__global__ void kernel(unsigned int jobs_todo, char* database, float* results ) {

  unsigned int id = threadIdx.x + blockIdx.x * blockDim.x;
  float A = 0; int i; char ch;
  if(id < jobs_todo) {
    for(i = 0; i < 1000; i += 1){
     ch = database[jobs_todo*i + id];
     if(ch == 'A') A++;
    }
    results[id] = A;
  }
}

И с char4 это

__global__ void kernel4(unsigned int jobs_todo, char4* database, float* results ) {

  unsigned int id = threadIdx.x + blockIdx.x * blockDim.x;
  float A = 0; int i; char4 ch4;
  if(id < jobs_todo) {
    for(i = 0; i < 1000/4; i += 1){
     ch4 = database[jobs_todo*i + id];
     if(ch4.x == 'A') A++;
     if(ch4.y == 'A') A++;
     if(ch4.z == 'A') A++;
     if(ch4.w == 'A') A++;
    }
    results[id] = A;
  }
}

Я также пробовал int4, но это всего на 0,0002 секунды быстрее, чем время char4.

person Ross    schedule 01.01.2013