Nvidia GPU 100 атомных транзакций - PullRequest
Новая
       6

Nvidia GPU 100 атомных транзакций

1 голос
/ 21 марта 2019

Я выполняю некоторые тесты атомарных (редукционных) транзакций одинарной точности, используя P100, и получаю случайные неожиданные результаты. Я надеюсь, что кто-то имеет представление о причине

Ниже приводится тестовая программа, которую я профилирую 

__global__ void atomic_test(float * out)
{
    int x=threadIdx.x+blockIdx.x*blockDim.x;
    //Create a pattern
    int myP=x/8;
    int myU=x%8;
    int myNewX=myP*myP*32+myU;
    atomicAdd(out+myNewX,1.0f);
}
main()
{
    int blocks=1;
    float * out;
    cudaMalloc(&out,sizeof(float)*16*32);
    dim3 threadDim;
    threadDim.x=32;
    threadDim.y=1;
    threadDim.z=1;
    for (int x=0;x<5;x++)
    atomic_test<<<blocks,threadDim>>>(out);
    cudaDeviceSynchronize();
    cudaFree(out);
    exit(0);
}

atomic_test запускается только с 1 деформацией, и все, что он делает, это атомарные добавления. Деформация каким-то образом разделена на 4, и каждая группа из 8 потоков выполнит атомарное добавление для правильно выровненного 32-битного слова.

Мое понимание P100 - любые транзакции, связанные с памятью, работают с 32-байтовыми выровненными словами, поэтому должно быть 4 атомарных транзакции, генерируемых Warp.

Странно то, что профилировщик много раз выдает 4 транзакции, как показано ниже 

 atomic_transactions               Atomic Transactions           4           4           4
 atomic_transactions_per_request   Atomic Transactions Per Request    4.000000    4.000000    4.000000
 l2_atomic_throughput              L2 Throughput (Atomic requests)  104.20MB/s  105.28MB/s  104.87MB/s
 l2_atomic_transactions            L2 Transactions (Atomic requests)          16          16          16

Но иногда сообщается о 6 транзакциях: 

 atomic_transactions               Atomic Transactions           6           6           6
 atomic_transactions_per_request   Atomic Transactions Per Request    6.000000    6.000000    6.000000
 l2_atomic_throughput              L2 Throughput (Atomic requests)  104.47MB/s  105.28MB/s  105.00MB/s
 l2_atomic_transactions            L2 Transactions (Atomic requests)          16          16          16

У кого-нибудь есть идея, почему это происходит? Может быть, есть проблемы с выравниванием? Кроме того, кто-нибудь знает, почему 4/6 атомарных транзакций создают 16 транзакций чтения на L2?

Заранее спасибо

...