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

Question

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

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

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

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

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

Answer 1

Я наконец понял ответ на свой вопрос. Определение с 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, но это просто на.0002 секунды быстрее, чем время char4.