Производительность процессора и памяти GPU - PullRequest
Новая
       43

Производительность процессора и памяти GPU

0 голосов
/ 29 июня 2019

Я хотел бы понять, как CUDA управляет доступом к ОЗУ GPU.Насколько мне известно, микросхемы ОЗУ ЦП и ОЗУ графического процессора не так уж отличаются, поэтому при справедливых условиях я ожидаю аналогичной производительности.

Я написал небольшой код, который суммирует элементы массива.Доступ к памяти может быть последовательным или разреженным, а массив может храниться в ОЗУ ЦП или ОЗУ графического процессора.Сумма выполняется одним потоком как в CPU, так и в GPU.

Код можно скомпилировать для строки комментариев CPU 7. 

#include <iostream>
#include <cassert>
#include <ctime>

using namespace std;

#define GPU true

#ifndef GPU
#define __global__
#define __host__
#define __device__
#endif

template<typename T>
__host__ 
void memAlloc(T** ptr, int count)
{
#ifdef GPU
    cudaMallocManaged(ptr, sizeof(T) * count);
#else
    *ptr = static_cast<T*>(malloc(sizeof(T) * count));
#endif
    assert(*ptr != nullptr);
}

__host__
int64_t getClockSpeed()
{
    int64_t clockSpeed;
#ifdef GPU
    int tmp;
    cudaDeviceGetAttribute(&tmp,cudaDevAttrClockRate,0);
    clockSpeed = static_cast<int64_t>(tmp) * 1000;
#else
    clockSpeed = static_cast<uint64_t>(CLOCKS_PER_SEC);
#endif    
    return clockSpeed / 1000;
}

static __device__ inline uint64_t gpuClock()
{
    uint64_t globaltime;
    asm volatile("mov.u64 %0, %%globaltimer;" : "=l"(globaltime) );
    return globaltime;
}

__device__
clock_t getTicks()
{
    clock_t ticks;
#ifdef GPU
    ticks = static_cast<clock_t>(gpuClock());
#else
    ticks = clock();
#endif
    return ticks;
}

__host__ __device__
int sparseRead(int32_t* mem, int elementCount, int stepSize, int rounds)
{
    int sum;

    for(int r = 0; r < rounds; r +=1)
    {
        sum = 0;
        for(int i = 0; i < stepSize; i += 1)
        {
            for(int j = i; j < elementCount; j += stepSize)
            {
                sum += mem[j];
            }
        }
    }

    return sum;
}

__host__ __device__
int sequentialRead(int32_t* mem, int elementCount, int rounds)
{
    int sum;

    for(int r = 0; r < rounds; r +=1)
    {
        sum = 0;
        for(int j = 0; j < elementCount; j += 1)
        {
            sum += mem[j];
        }   
    }

    return sum;
}

__global__
void getElapsedTicks(int32_t* mem, int elementCount, int stepSize, int rounds, clock_t* elapsedTicks, int32_t* sum)
{
    for(int i = 0; i < elementCount; i +=1)
    {
        mem[i] = i % rounds;
    }

    clock_t start = getTicks();
    printf("Start %d\n", start);

    *sum = sparseRead(mem, elementCount, stepSize, rounds);
    //*sum = sequentialRead(mem, elementCount, rounds);

    clock_t end =  getTicks();
    printf("End %d\n", end);

    *elapsedTicks = end - start;
}

int main(int argc, char *argv[]) 
{
    int elementCount = atoi(argv[1]);
    int stepSize = atoi(argv[2]);
    int rounds = atoi(argv[3]);

    int32_t* mem;
    memAlloc(&mem, elementCount);

    clock_t* elapsedTicks;
    memAlloc(&elapsedTicks, 1);

    int32_t* sum;
    memAlloc(&sum, 1);

#ifdef GPU
    getElapsedTicks<<<1,1>>>(mem, elementCount, stepSize, rounds, elapsedTicks, sum); 
    cudaDeviceSynchronize();
#else
    getElapsedTicks(mem, elementCount, stepSize, rounds, elapsedTicks, sum); 
#endif

    uint64_t elapsedTime = *elapsedTicks / getClockSpeed();
    cout << "Sum "<< *sum << endl;
    cout << "Elapsed " << elapsedTime << " ms" << endl;
}

Я ожидал ~ 3-кратное замедление на GPU из-за более медленной тактовой частотыно у меня есть 100-кратное замедление для последовательных и разреженных сумм.Чего мне не хватает?

...