Cuda: один битовый набор XOR с массивом битовых наборов - PullRequest
Новая
       106

Cuda: один битовый набор XOR с массивом битовых наборов

0 голосов
/ 24 января 2020

Я хочу XOR одного набора битов с кучей других наборов битов (~ 100k) и подсчитать установленные биты каждого результата xor. Размер одного набора битов составляет около 20 000 бит.

Наборы битов уже преобразованы в массивы unsigned int, чтобы можно было использовать функцию intrinsi c __popc(). «Связка» уже находится в памяти устройства.

Мой текущий код ядра выглядит следующим образом: 

// Grid/Blocks used for kernel invocation 
dim3 block(32); 
dim3 grid((bunch_size / 31) + 32);

__global__ void kernelXOR(uint * bitset, uint * bunch, int * set_bits, int bitset_size, int bunch_size) {

    int tid = blockIdx.x*blockDim.x + threadIdx.x;

    if (tid < bunch_size){      // 1 Thread for each bitset in the 'bunch'
        int sum = 0;
        uint xor_res = 0;
        for (int i = 0; i < bitset_size; ++i){  // Iterate through every uint-block of the bitsets
            xor_res = bitset[i] ^ bunch[bitset_size * tid + i];
            sum += __popc(xor_res);
        }
        set_bits[tid] = sum;
    }
}

Однако по сравнению с распараллеленной версией c ++ / boost я не вижу Преимущество использования Cuda.

Есть ли потенциал для оптимизации этого ядра?

1 Ответ

3 голосов
/ 24 января 2020

Есть ли потенциал для оптимизации этого ядра?

Здесь я вижу 2 проблемы (и они являются первыми двумя классическими первичными целями оптимизации для любого программиста CUDA):

  1. Вы хотите попытаться эффективно использовать глобальную память. Ваш доступ к bitset и bunch не объединен. (эффективно используйте подсистемы памяти)

  2. Использование 32 потоков на блок обычно не рекомендуется и может ограничить общую занятость. Один поток на битовый набор также потенциально проблематичен c. (раскрыть достаточно параллелизма)

Будет ли решение этих проблем соответствовать вашему определению выгоды, невозможно определить без контрольного теста. Кроме того, простые связанные с памятью проблемы, подобные этой, редко интересны в CUDA, когда рассматриваются сами по себе. Тем не менее, мы можем (вероятно) улучшить производительность вашего ядра.

Мы будем использовать отобранный dry список идей:

  • каждый блок обрабатывает набор битов, скорее чем каждый поток, чтобы включить объединение
  • использовать общую память для загрузки набора битов сравнения и использовать его повторно
  • использовать достаточно блоков для насыщения графического процессора вместе с шагами шага
  • используйте украшение в стиле const ... __restrict__, чтобы извлечь выгоду из кэша RO

Вот рабочий пример: 

$ cat t1649.cu
#include <iostream>
#include <cstdlib>

const int my_bitset_size = 20000/(32);
const int my_bunch_size = 100000;
typedef unsigned uint;

//using one thread per bitset in the bunch
__global__ void kernelXOR(uint * bitset, uint * bunch, int * set_bits, int bitset_size, int bunch_size) {

    int tid = blockIdx.x*blockDim.x + threadIdx.x;

    if (tid < bunch_size){      // 1 Thread for each bitset in the 'bunch'
        int sum = 0;
        uint xor_res = 0;
        for (int i = 0; i < bitset_size; ++i){  // Iterate through every uint-block of the bitsets
            xor_res = bitset[i] ^ bunch[bitset_size * tid + i];
            sum += __popc(xor_res);
        }
        set_bits[tid] = sum;
    }
}

const int nTPB = 256;
// one block per bitset, multiple bitsets per block
__global__ void kernelXOR_imp(const uint * __restrict__  bitset, const uint * __restrict__  bunch, int * __restrict__  set_bits, int bitset_size, int bunch_size) {

    __shared__ uint sbitset[my_bitset_size];  // could also be dynamically allocated for varying bitset sizes
    __shared__ int ssum[nTPB];
    // load shared, block-stride loop
    for (int idx = threadIdx.x; idx < bitset_size; idx += blockDim.x) sbitset[idx] = bitset[idx];
    __syncthreads();
    // stride across all bitsets in bunch
    for (int bidx = blockIdx.x; bidx < bunch_size; bidx += gridDim.x){
      int my_sum = 0;
      for (int idx = threadIdx.x; idx < bitset_size; idx += blockDim.x) my_sum += __popc(sbitset[idx] ^ bunch[bidx*bitset_size + idx]);
    // block level parallel reduction
      ssum[threadIdx.x] = my_sum;
      for (int ridx = nTPB>>1; ridx > 0; ridx >>=1){
        __syncthreads();
        if (threadIdx.x < ridx) ssum[threadIdx.x] += ssum[threadIdx.x+ridx];}
      if (!threadIdx.x) set_bits[bidx] = ssum[0];}
}



int main(){

// data setup

  uint *d_cbitset, *d_bitsets, *h_cbitset, *h_bitsets;
  int *d_r, *h_r, *h_ri;
  h_cbitset = new uint[my_bitset_size];
  h_bitsets = new uint[my_bitset_size*my_bunch_size];
  h_r  = new int[my_bunch_size];
  h_ri = new int[my_bunch_size];
  for (int i = 0; i < my_bitset_size*my_bunch_size; i++){
    h_bitsets[i] = rand();
    if (i < my_bitset_size) h_cbitset[i] = rand();}
  cudaMalloc(&d_cbitset, my_bitset_size*sizeof(uint));
  cudaMalloc(&d_bitsets, my_bitset_size*my_bunch_size*sizeof(uint));
  cudaMalloc(&d_r,  my_bunch_size*sizeof(int));
  cudaMemcpy(d_cbitset, h_cbitset, my_bitset_size*sizeof(uint), cudaMemcpyHostToDevice);
  cudaMemcpy(d_bitsets, h_bitsets, my_bitset_size*my_bunch_size*sizeof(uint), cudaMemcpyHostToDevice);
// original

// Grid/Blocks used for kernel invocation
  dim3 block(32);
  dim3 grid((my_bunch_size / 31) + 32);

  kernelXOR<<<grid, block>>>(d_cbitset, d_bitsets, d_r, my_bitset_size, my_bunch_size);
  cudaMemcpy(h_r, d_r, my_bunch_size*sizeof(int), cudaMemcpyDeviceToHost);


// improved
  dim3 iblock(nTPB);
  dim3 igrid(640);
  kernelXOR_imp<<<igrid, iblock>>>(d_cbitset, d_bitsets, d_r, my_bitset_size, my_bunch_size);
  cudaMemcpy(h_ri, d_r, my_bunch_size*sizeof(int), cudaMemcpyDeviceToHost);

  for (int i = 0; i < my_bunch_size; i++)
    if (h_r[i] != h_ri[i]) {std::cout << "mismatch at i: " << i << " was: " << h_ri[i] << " should be: " << h_r[i] << std::endl; return 0;}
  std::cout << "Results match." << std::endl;
  return 0;
}
$ nvcc -o t1649 t1649.cu
$ cuda-memcheck ./t1649
========= CUDA-MEMCHECK
Results match.
========= ERROR SUMMARY: 0 errors
$ nvprof ./t1649
==18868== NVPROF is profiling process 18868, command: ./t1649
Results match.
==18868== Profiling application: ./t1649
==18868== Profiling result:
            Type  Time(%)      Time     Calls       Avg       Min       Max  Name
 GPU activities:   97.06%  71.113ms         2  35.557ms  2.3040us  71.111ms  [CUDA memcpy HtoD]
                    2.26%  1.6563ms         1  1.6563ms  1.6563ms  1.6563ms  kernelXOR(unsigned int*, unsigned int*, int*, int, int)
                    0.59%  432.68us         1  432.68us  432.68us  432.68us  kernelXOR_imp(unsigned int const *, unsigned int const *, int*, int, int)
                    0.09%  64.770us         2  32.385us  31.873us  32.897us  [CUDA memcpy DtoH]
      API calls:   78.20%  305.44ms         3  101.81ms  11.373us  304.85ms  cudaMalloc
                   18.99%  74.161ms         4  18.540ms  31.554us  71.403ms  cudaMemcpy
                    1.39%  5.4121ms         4  1.3530ms  675.30us  3.3410ms  cuDeviceTotalMem
                    1.26%  4.9393ms       388  12.730us     303ns  530.95us  cuDeviceGetAttribute
                    0.11%  442.37us         4  110.59us  102.61us  125.59us  cuDeviceGetName
                    0.03%  128.18us         2  64.088us  21.789us  106.39us  cudaLaunchKernel
                    0.01%  35.764us         4  8.9410us  2.9670us  18.982us  cuDeviceGetPCIBusId
                    0.00%  8.3090us         8  1.0380us     540ns  1.3870us  cuDeviceGet
                    0.00%  5.9530us         3  1.9840us     310ns  3.9900us  cuDeviceGetCount
                    0.00%  2.8800us         4     720ns     574ns     960ns  cuDeviceGetUuid
$

В этом случае на моем Tesla V100 Что касается размера вашей проблемы, я вижу увеличение производительности ядра в 4 раза. Однако производительность ядра здесь ничтожна по сравнению со стоимостью перемещения данных. Поэтому маловероятно, что такого рода оптимизации будут иметь существенное значение в вашем сравнительном тестовом случае, если это единственное, что вы делаете на графическом процессоре.

Приведенный выше код использует зацикливание на уровне блоков и на уровне сетки, что означает, что он должен вести себя корректно практически для любого размера блока потоков (кратного 32, пожалуйста), а также размера сетки. Это не значит, что любой / все варианты будут одинаковыми. Выбор размера ниточного блока должен обеспечить возможность почти полного заполнения (поэтому не выбирайте 32). Выбор размера сетки - это количество блоков для достижения полной занятости на SM, умноженное на количество SM. Это должны быть почти оптимальные варианты, но, согласно моим испытаниям, например, большее количество блоков на самом деле не снижает производительность, а производительность должна быть примерно постоянной для почти любого размера потокового блока (кроме 32), при условии, что количество блоков рассчитывается соответствующим образом. .

...