Атомный счетчик для критической секции, не использующий атомную полосу пропускания согласно профилировщику - PullRequest
Новая
       39

Атомный счетчик для критической секции, не использующий атомную полосу пропускания согласно профилировщику

0 голосов
/ 15 ноября 2018

Несмотря на то, что для общей переменной есть первичный доступ для общей переменной для каждого потока, профилировщик показывает нулевую пропускную способность для атомарных элементов:

enter image description here

Минимальное воспроизведениеПример, который я мог бы сделать здесь: 

#include <stdio.h>
#include <cuda_runtime.h>

#define criticalSection(T, ...) {\
    __shared__ int ctrBlock; \
    if(threadIdx.x==0) \
       ctrBlock=0; \
    __syncthreads(); \
    while(atomicAdd(&ctrBlock,0)<(blockDim.x/32)) \
    { \
       if( atomicAdd(&ctrBlock,0) == (threadIdx.x/32) ) \
       { \
            int ctr=0; \
            while(ctr<32) \
            { \
                   if( ctr == (threadIdx.x&31) ) \
                   { \
                    { \
                          T,##__VA_ARGS__; \
                    } \
                   } \
                   ctr++; \
                   __syncwarp(); \
            } \
            if((threadIdx.x&31) == 0)atomicAdd(&ctrBlock,1); \
        } \
        __syncthreads(); \
     } \
}

__global__ void
vectorAdd(const float *A, const float *B, float *C, int numElements)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;

     // instead of if(i==0) C[0]=0.0f; initialization
    if(i==blockDim.x*blockIdx.x)
       C[blockDim.x*blockIdx.x]=0.0f;

    __syncthreads();
    criticalSection({
        if (i < numElements)
        {
           C[blockDim.x*blockIdx.x] += A[i] + B[i];
        }
    });
}


int main(void)
{
    int numElements = 50000;
    size_t size = numElements * sizeof(float);
    float *h_A = (float *)malloc(size); 
    float *h_B = (float *)malloc(size);
    float *h_C = (float *)malloc(size);

    for (int i = 0; i < numElements; ++i)
    {
        h_A[i] = i;
        h_B[i] = 2*i;
    }

    float *d_A = NULL;
    cudaMalloc((void **)&d_A, size);

    float *d_B = NULL;
    cudaMalloc((void **)&d_B, size);

    float *d_C = NULL;
    cudaMalloc((void **)&d_C, size);

    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
    int threadsPerBlock = 256;
    int blocksPerGrid =(numElements + threadsPerBlock - 1) / threadsPerBlock;
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements);
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    printf("%g\n",h_C[0]);


    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    free(h_A);
    free(h_B);
    free(h_C);

    return 0;
}

, он правильно выводит сумму (от 1 до 255) * 3 результата (на каждый начальный элемент в блоке) при каждом запуске.

Вопрос: почему профилировщик показывает, что он не использует атомную полосу пропускания, даже если он правильно работает?

Ядро завершает (196 блоков, 256 потоков на блок) менее 2,4 миллисекунд на 192-ядерном KeplerGPU.Собирает ли GPU атомарные элементы и преобразовывает их во что-то более эффективное в каждой точке синхронизации?

Это не дает никакой ошибки, я удалил проверку ошибок на удобочитаемость.

Изменение добавления элемента массива C в:

((volatile float *) C)[blockDim.x*blockIdx.x] += A[i] + B[i];

не меняет ни поведение, ни результат.

Использование инструментария CUDA 9.2 и драйвера v396, Ubuntu 16.04, Quadro K420.

Вот команды компиляции: 

nvcc -ccbin g++ -m64 -gencode arch=compute_30,code=sm_30 -gencode arch=compute_35,code=sm_35 -gencode arch=compute_37,code=sm_37 -gencode arch=compute_50,code=sm_50 -gencode arch=compute_52,code=sm_52 -gencode arch=compute_60,code=sm_60 -gencode arch=compute_61,code=sm_61 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_70,code=compute_70 -o vectorAdd.o -c vectorAdd.cu
nvcc -ccbin g++ -m64 -gencode arch=compute_30,code=sm_30 -gencode arch=compute_35,code=sm_35 -gencode arch=compute_37,code=sm_37 -gencode arch=compute_50,code=sm_50 -gencode arch=compute_52,code=sm_52 -gencode arch=compute_60,code=sm_60 -gencode arch=compute_61,code=sm_61 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_70,code=compute_70 -o vectorAdd vectorAdd.o

Ptx-вывод cuobjdump (sass было более 50 тыс. Символов): 

.visible .entry _Z9vectorAddPKfS0_Pfi(
.param .u64 _Z9vectorAddPKfS0_Pfi_param_0,
.param .u64 _Z9vectorAddPKfS0_Pfi_param_1,
.param .u64 _Z9vectorAddPKfS0_Pfi_param_2,
.param .u32 _Z9vectorAddPKfS0_Pfi_param_3
)
{
.reg .pred %p<32>;
.reg .f32 %f<41>;
.reg .b32 %r<35>;
.reg .b64 %rd<12>;

    .shared .align 4 .u32 _ZZ9vectorAddPKfS0_PfiE8ctrBlock;

ld.param.u64 %rd5, [_Z9vectorAddPKfS0_Pfi_param_0];
ld.param.u64 %rd6, [_Z9vectorAddPKfS0_Pfi_param_1];
ld.param.u64 %rd7, [_Z9vectorAddPKfS0_Pfi_param_2];
ld.param.u32 %r13, [_Z9vectorAddPKfS0_Pfi_param_3];
cvta.to.global.u64 %rd1, %rd7;
mov.u32 %r14, %ctaid.x;
mov.u32 %r1, %ntid.x;
mul.lo.s32 %r2, %r14, %r1;
mov.u32 %r3, %tid.x;
add.s32 %r4, %r2, %r3;
setp.ne.s32 %p8, %r4, 0;
@%p8 bra BB0_2;

mov.u32 %r15, 0;
st.global.u32 [%rd1], %r15;

BB0_2:
bar.sync 0;
setp.ne.s32 %p9, %r3, 0;
@%p9 bra BB0_4;

mov.u32 %r16, 0;
st.shared.u32 [_ZZ9vectorAddPKfS0_PfiE8ctrBlock], %r16;

BB0_4:
bar.sync 0;
mov.u32 %r17, _ZZ9vectorAddPKfS0_PfiE8ctrBlock;
atom.shared.add.u32 %r18, [%r17], 0;
shr.u32 %r5, %r1, 5;
setp.ge.u32 %p10, %r18, %r5;
@%p10 bra BB0_27;

shr.u32 %r6, %r3, 5;
and.b32 %r7, %r3, 31;
cvta.to.global.u64 %rd8, %rd5;
mul.wide.s32 %rd9, %r4, 4;
add.s64 %rd2, %rd8, %rd9;
cvta.to.global.u64 %rd10, %rd6;
add.s64 %rd3, %rd10, %rd9;
mul.wide.u32 %rd11, %r2, 4;
add.s64 %rd4, %rd1, %rd11;
neg.s32 %r8, %r7;

BB0_6:
atom.shared.add.u32 %r21, [%r17], 0;
mov.u32 %r34, 0;
setp.ne.s32 %p11, %r21, %r6;
mov.u32 %r33, %r8;
@%p11 bra BB0_26;

BB0_7:
setp.eq.s32 %p12, %r33, 0;
setp.lt.s32 %p13, %r4, %r13;
and.pred %p14, %p12, %p13;
@!%p14 bra BB0_9;
bra.uni BB0_8;

BB0_8:
ld.global.f32 %f1, [%rd2];
ld.global.f32 %f2, [%rd3];
add.f32 %f3, %f1, %f2;
ld.volatile.global.f32 %f4, [%rd4];
add.f32 %f5, %f4, %f3;
st.volatile.global.f32 [%rd4], %f5;

BB0_9:
bar.warp.sync -1;
add.s32 %r22, %r34, 1;
setp.eq.s32 %p15, %r22, %r7;
and.pred %p16, %p15, %p13;
@!%p16 bra BB0_11;
bra.uni BB0_10;

BB0_10:
ld.global.f32 %f6, [%rd2];
ld.global.f32 %f7, [%rd3];
add.f32 %f8, %f6, %f7;
ld.volatile.global.f32 %f9, [%rd4];
add.f32 %f10, %f9, %f8;
st.volatile.global.f32 [%rd4], %f10;

BB0_11:
bar.warp.sync -1;
add.s32 %r23, %r34, 2;
setp.eq.s32 %p17, %r23, %r7;
and.pred %p18, %p17, %p13;
@!%p18 bra BB0_13;
bra.uni BB0_12;

BB0_12:
ld.global.f32 %f11, [%rd2];
ld.global.f32 %f12, [%rd3];
add.f32 %f13, %f11, %f12;
ld.volatile.global.f32 %f14, [%rd4];
add.f32 %f15, %f14, %f13;
st.volatile.global.f32 [%rd4], %f15;

BB0_13:
bar.warp.sync -1;
add.s32 %r24, %r34, 3;
setp.eq.s32 %p19, %r24, %r7;
and.pred %p20, %p19, %p13;
@!%p20 bra BB0_15;
bra.uni BB0_14;

BB0_14:
ld.global.f32 %f16, [%rd2];
ld.global.f32 %f17, [%rd3];
add.f32 %f18, %f16, %f17;
ld.volatile.global.f32 %f19, [%rd4];
add.f32 %f20, %f19, %f18;
st.volatile.global.f32 [%rd4], %f20;

BB0_15:
bar.warp.sync -1;
add.s32 %r25, %r34, 4;
setp.eq.s32 %p21, %r25, %r7;
and.pred %p22, %p21, %p13;
@!%p22 bra BB0_17;
bra.uni BB0_16;

BB0_16:
ld.global.f32 %f21, [%rd2];
ld.global.f32 %f22, [%rd3];
add.f32 %f23, %f21, %f22;
ld.volatile.global.f32 %f24, [%rd4];
add.f32 %f25, %f24, %f23;
st.volatile.global.f32 [%rd4], %f25;

BB0_17:
bar.warp.sync -1;
add.s32 %r26, %r34, 5;
setp.eq.s32 %p23, %r26, %r7;
and.pred %p24, %p23, %p13;
@!%p24 bra BB0_19;
bra.uni BB0_18;

BB0_18:
ld.global.f32 %f26, [%rd2];
ld.global.f32 %f27, [%rd3];
add.f32 %f28, %f26, %f27;
ld.volatile.global.f32 %f29, [%rd4];
add.f32 %f30, %f29, %f28;
st.volatile.global.f32 [%rd4], %f30;

BB0_19:
bar.warp.sync -1;
add.s32 %r27, %r34, 6;
setp.eq.s32 %p25, %r27, %r7;
and.pred %p26, %p25, %p13;
@!%p26 bra BB0_21;
bra.uni BB0_20;

BB0_20:
ld.global.f32 %f31, [%rd2];
ld.global.f32 %f32, [%rd3];
add.f32 %f33, %f31, %f32;
ld.volatile.global.f32 %f34, [%rd4];
add.f32 %f35, %f34, %f33;
st.volatile.global.f32 [%rd4], %f35;

BB0_21:
bar.warp.sync -1;
add.s32 %r28, %r34, 7;
setp.eq.s32 %p27, %r28, %r7;
and.pred %p28, %p27, %p13;
@!%p28 bra BB0_23;
bra.uni BB0_22;

BB0_22:
ld.global.f32 %f36, [%rd2];
ld.global.f32 %f37, [%rd3];
add.f32 %f38, %f36, %f37;
ld.volatile.global.f32 %f39, [%rd4];
add.f32 %f40, %f39, %f38;
st.volatile.global.f32 [%rd4], %f40;

BB0_23:
add.s32 %r34, %r34, 8;
bar.warp.sync -1;
add.s32 %r33, %r33, 8;
setp.ne.s32 %p29, %r34, 32;
@%p29 bra BB0_7;

setp.ne.s32 %p30, %r7, 0;
@%p30 bra BB0_26;

atom.shared.add.u32 %r30, [%r17], 1;

BB0_26:
bar.sync 0;
atom.shared.add.u32 %r32, [%r17], 0;
setp.lt.u32 %p31, %r32, %r5;
@%p31 bra BB0_6;

BB0_27:
ret;
}

1 Ответ

0 голосов
/ 16 ноября 2018

Здесь нужно знать как минимум 2 вещи.

  1. Заметим, что ваша программа использует атомарные элементы в общей памяти местах.Кроме того, вы указали, что вы компилируете (и при профилировании запускаете) графический процессор с архитектурой Kepler.

    В Kepler атомы общей памяти эмулируются с помощью программной последовательности .Это не будет видно при проверке кода PTX, поскольку преобразование в последовательность эмуляции выполняется ptxas, инструментом, который преобразует PTX в код SASS для выполнения на целевом устройстве.

    Поскольку вынацеленный и работающий на Kepler, SASS не содержит атомарных инструкций для совместно используемой памяти (вместо этого совместно используемые атомы эмулируются с помощью цикла, который использует специальные аппаратные блокировки, и, например, вы можете увидеть LDSLK, инструкцию загрузки из общего доступа с блокировкой,в вашем коде SASS).

    Поскольку в вашем коде нет фактических атомарных инструкций (в Kepler), он не генерирует атомарный трафик, который может быть отслежен профилировщиком.

    Если вы хотите проверитьдля этого используйте cuobjdump tool в скомпилированном двоичном файле.Я рекомендую компилировать only для целевой архитектуры Kepler, которую вы фактически будете использовать для этого вида бинарного анализа.Вот пример: 

    $ nvcc -o t324 t324.cu -arch=sm_30
$ cuobjdump -sass ./t324 |grep ATOM
$ nvcc -o t324 t324.cu -arch=sm_50
$ cuobjdump -sass ./t324 |grep ATOM
        /*00e8*/               @P2 ATOMS.ADD R6, [RZ], RZ ;                       /* 0xec0000000ff2ff06 */
        /*01b8*/               @P0 ATOMS.ADD R12, [RZ], RZ ;                      /* 0xec0000000ff0ff0c */
        /*10f8*/               @P0 ATOMS.ADD RZ, [RZ], R12 ;                      /* 0xec00000000c0ffff */
        /*1138*/               @P0 ATOMS.ADD R10, [RZ], RZ ;                      /* 0xec0000000ff0ff0a */
$

  2. Как указано выше, для Максвелла и далее существует собственная атомарная инструкция общей памяти (например, ATOMS) в коде SASS.Поэтому, если вы скомпилируете свой код для архитектуры maxwell или выше, вы увидите фактические атомарные инструкции в SASS.

    Однако я не уверен, будет ли это отображаться в визуальном профилировщике или как это будет.Я подозреваю, что общая атомная отчетность может быть ограничена.Это можно обнаружить, просмотрев доступные метрики и заметив, что для архитектур 5.0 и выше большинство атомных метрик предназначено специально для глобальных атомик, и единственная метрика, которую я могу найти относительно общей атомности: 

    inst_executed_shared_atomics    Warp level shared instructions for atom and atom CAS    Multi-context

    Я не уверен, что этого достаточно для вычисления пропускной способности или использования, поэтому я не уверен, что визуальный профилировщик намеревается сообщить много о способе совместного использования атома, даже на архитектурах 5.0+.Вы можете попробовать это, конечно.

В качестве отступления, я обычно думаю, что такая конструкция подразумевает логический дефект в коде: 

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

if(i==0)
   C[0]=0.0f;
__syncthreads();

Но это не относится к данному конкретному запросу, и я все равно не уверен в намерении вашего кода.Помните, что в CUDA не указан порядок выполнения блока.

...