Какая альтернатива для __match_any_syn c в вычислительных возможностях 6? - PullRequest
Новая
       18

Какая альтернатива для __match_any_syn c в вычислительных возможностях 6?

2 голосов
/ 23 января 2020

В примерах cuda используется , например, здесь , __match_all_sync __match_any_sync.

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

// increment the value at ptr by 1 and return the old value
__device__ int atomicAggInc(int* ptr) {
    int pred;
    //const auto mask = __match_all_sync(__activemask(), ptr, &pred); //error, should be any_sync, not all_sync
    const auto mask = __match_any_sync(__activemask(), ptr, &pred);
    const auto leader = __ffs(mask) - 1;  // select a leader
    int res;
    const auto lane_id = ThreadId() % warpSize;
    if (lane_id == leader) {                 // leader does the update
        res = atomicAdd(ptr, __popc(mask));
    }
    res = __shfl_sync(mask, res, leader);    // get leader’s old value
    return res + __popc(mask & ((1 << lane_id) - 1)); //compute old value
}

Здесь __match_any_sync разделяет потоки в деформации на группы, имеющие одинаковое значение ptr, чтобы каждая группа могла обновлять свой собственный ptr атомарно, не мешая другим потокам.

Я знаю, что компилятор nv cc (начиная с cuda 9) автоматически выполняет такую ​​оптимизацию под капотом, но это всего лишь механика __match_any_sync

Есть ли способ сделать эту возможность до вычисления 7?

1 Ответ

4 голосов
/ 23 января 2020

РЕДАКТИРОВАТЬ: Теперь статья блога была изменена, чтобы отразить __match_any_sync(), а не __match_all_sync(), поэтому любые комментарии на этот счет следует игнорировать. Ответ ниже отредактирован, чтобы отразить это.

Основываясь на вашем утверждении:

это как раз о механике __match_any_sync

мы будем сосредоточиться на замене самого __match_any_sync, а не на какой-либо другой форме переписывания функции atomicAggInc. Поэтому мы должны предоставить маску с тем же значением, которое будет возвращено __match_any_sync() на архитектурах cc7.0 или выше.

Я полагаю, что для этого потребуется al oop, который транслирует ptr значение, в худшем случае одна итерация для каждого потока в деформации (поскольку каждый поток может иметь уникальное значение ptr) и проверка того, какие потоки имеют одинаковое значение. Существуют различные способы, которыми мы могли бы «оптимизировать» эту l oop для этой функции, чтобы, возможно, уменьшить количество отключений с 32 до некоторого меньшего значения, основываясь на фактических значениях ptr в каждом потоке, но такая оптимизация в моем представление вносит значительную сложность, что увеличивает время обработки в худшем случае (как это типично для оптимизаций с ранним выходом). Поэтому я продемонстрирую довольно простой метод без этой оптимизации.

Другое соображение - что делать в случае, если основы не сходятся? Для этого мы можем использовать __activemask() для идентификации этого случая.

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

$ cat t1646.cu
#include <iostream>
#include <stdio.h>

// increment the value at ptr by 1 and return the old value
__device__ int atomicAggInc(int* ptr) {
    int mask;
#if __CUDA_ARCH__ >= 700
    mask = __match_any_sync(__activemask(), (unsigned long long)ptr);
#else
    unsigned tmask = __activemask();
    for (int i = 0; i < warpSize; i++){
#ifdef USE_OPT
      if ((1U<<i) & tmask){
#endif
        unsigned long long tptr = __shfl_sync(tmask, (unsigned long long)ptr, i);
        unsigned my_mask = __ballot_sync(tmask, (tptr == (unsigned long long)ptr));
        if (i == (threadIdx.x & (warpSize-1))) mask = my_mask;}
#ifdef USE_OPT
      }
#endif
#endif
    int leader = __ffs(mask) - 1;  // select a leader
    int res;
    unsigned lane_id = threadIdx.x % warpSize;
    if (lane_id == leader) {                 // leader does the update
        res = atomicAdd(ptr, __popc(mask));
    }
    res = __shfl_sync(mask, res, leader);    // get leader’s old value
    return res + __popc(mask & ((1 << lane_id) - 1)); //compute old value
}



__global__ void k(int *d){

  int *ptr = d + threadIdx.x/4;
  if ((threadIdx.x >= 16) && (threadIdx.x < 32))
    atomicAggInc(ptr);
}

const int ds = 32;
int main(){

  int *d_d, *h_d;
  h_d = new int[ds];
  cudaMalloc(&d_d, ds*sizeof(d_d[0]));
  cudaMemset(d_d, 0, ds*sizeof(d_d[0]));
  k<<<1,ds>>>(d_d);
  cudaMemcpy(h_d, d_d, ds*sizeof(d_d[0]), cudaMemcpyDeviceToHost);
  for (int i = 0; i < ds; i++)
    std::cout << h_d[i] << " ";
  std::cout << std::endl;
}
$ nvcc -o t1646 t1646.cu -DUSE_OPT
$ cuda-memcheck ./t1646
========= CUDA-MEMCHECK
0 0 0 0 4 4 4 4 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
========= ERROR SUMMARY: 0 errors
$

(CentOS 7, CUDA 10.1.243, где устройство 0 - Tesla V100, устройство 1 - устройство cc3.5).

Я добавил дополнительную оптимизацию для случая, когда деформация отклонена (т. Е. tmask не 0xFFFFFFFF). Это можно выбрать, определив USE_OPT.

...