Оптимизация использования конкретной памяти для CUDA

Question

У меня есть задача обработки данных, которую можно стилизовать следующим образом. я имею data (~1-10GB) и функция, которая генерирует summary (~1MB) на основании этого data и некоторый (двойной) ввод x, Мне нужно получить это summary для ~1000 значений xчто выглядело как идеальное задание для GPU. Повторим, вход data одинаково для всех потоков и читается линейно, но каждый поток должен создавать свой собственный summary, Функции выполняются независимо для разных x,

Тем не менее, грубый однопоточный цикл через все значения x на CPU дает только в 3 раза худшую производительность, чем K520. Я понимаю, что это задача, требующая большого объема памяти (потоки должны обращаться к случайным частям и записывать их summary), но я все еще пытаюсь понять, как GPU может потерять свое первоначальное преимущество в 1000 раз. Я пытался кормить data кормить кусками используя __constant__ память (так как это один и тот же ввод во всех потоках), без видимых улучшений. Типичное время выполнения блока, как сообщает nvprof, составляет 10-30 секунд.

Буду признателен за понимание оптимизации, подходящей для этой задачи.

РЕДАКТИРОВАТЬ: Ниже приведен пример кода, который повторяет проблему. Его можно скомпилировать как в g++ (время выполнения отчета 5 с), так и в nvcc (время выполнения отчета 7 с). Результаты профилирования следующие

==23844== Результат профилирования:
Время (%) Время звонков Средняя Мин Макс Имя
98.86% 4.68899s 1 4.68899s 4.68899s 4.68899s Ядро (Наблюдение *, int*, Информация **)
1,09% 51,480мс 4 12,870мс 1,9200ус 50,426мс [CUDA memcpy HtoD]
0,06% 2,6634 мс 800 3,390 мкс 3,2950 мкс 5,100 мкс [CUDA memcpy DtoD]
0.00% 4.3200us 1 4.3200us 4.3200us 4.3200us [CUDA memcpy DtoH]

#include <iostream>
#include <fstream>
#include <cstdlib>
#include <ctime>
#include <cstring>

#define MAX_OBS 1000000
#define MAX_BUCKETS 1000

using namespace std;

// Cross-arch defines
#ifndef __CUDACC__

#define GPU_FUNCTION

#define cudaSuccess 0

typedef int cudaError_t;

struct dim3
{
    int x;
    int y;
    int z;
} blockIdx, threadIdx;

enum cudaMemcpyKind
{
    cudaMemcpyHostToDevice = 0,
    cudaMemcpyDeviceToHost = 1, 
    cudaMemcpyDeviceToDevice = 2
};

cudaError_t cudaMalloc(void ** Dst, size_t bytes)
{
    return !(*Dst = malloc(bytes));
}

cudaError_t cudaMemcpy(void * Dst, const void * Src, size_t bytes, cudaMemcpyKind kind)
{
    return !memcpy(Dst, Src, bytes);
}

#else
#define GPU_FUNCTION __global__
#endif

// Basic observation structure as stored on disk
struct Observation
{
    double core[20];
};

struct Info
{
    int left;
    int right;
};

GPU_FUNCTION void Kernel(Observation * d_obs, 
                         int * d_bucket,
                         Info ** d_summaries)
{
    Info * summary = d_summaries[threadIdx.x * 40 + threadIdx.y];

    for (int i = 0; i < MAX_OBS; i++)
    {
        if (d_obs[i].core[threadIdx.x] < (threadIdx.x + 1) * threadIdx.y)
            summary[d_bucket[i]].left++;
        else
            summary[d_bucket[i]].right++;
    }
}

int main()
{
    srand((unsigned int)time(NULL));

    // Generate dummy observations
    Observation * obs = new Observation [MAX_OBS];
    for (int i = 0; i < MAX_OBS; i++)
        for (int j = 0; j < 20; j++)
            obs[i].core[j] = (double)rand() / RAND_MAX;

    // Attribute observations to one of the buckets
    int * bucket = new int [MAX_OBS];
    for (int i = 0; i < MAX_OBS; i++)
        bucket[i] = rand() % MAX_BUCKETS;

    Info summary[MAX_BUCKETS];
    for (int i = 0; i < MAX_BUCKETS; i++)
        summary[i].left = summary[i].right = 0;

    time_t start;
    time(&start);

    // Init device objects
    Observation * d_obs;                    
    int * d_bucket; 
    Info * d_summary;
    Info ** d_summaries;

    cudaMalloc((void**)&d_obs, MAX_OBS * sizeof(Observation));
    cudaMemcpy(d_obs, obs, MAX_OBS * sizeof(Observation), cudaMemcpyHostToDevice);
    cudaMalloc((void**)&d_bucket, MAX_OBS * sizeof(int));
    cudaMemcpy(d_bucket, bucket, MAX_OBS * sizeof(int), cudaMemcpyHostToDevice);
    cudaMalloc((void**)&d_summary, MAX_BUCKETS * sizeof(Info));
    cudaMemcpy(d_summary, summary, MAX_BUCKETS * sizeof(Info), cudaMemcpyHostToDevice);

    Info ** tmp_summaries = new Info * [20 * 40];
    for (int k = 0; k < 20 * 40; k++)           
        cudaMalloc((void**)&tmp_summaries[k], MAX_BUCKETS * sizeof(Info));
    cudaMalloc((void**)&d_summaries, 20 * 40 * sizeof(Info*));
    cudaMemcpy(d_summaries, tmp_summaries, 20 * 40 * sizeof(Info*), cudaMemcpyHostToDevice);
    for (int k = 0; k < 20 * 40; k++)
        cudaMemcpy(tmp_summaries[k], d_summary, MAX_BUCKETS * sizeof(Info), cudaMemcpyDeviceToDevice);

#ifdef __CUDACC__
    Kernel<<<1, dim3(20, 40, 1)>>>(d_obs, d_bucket, d_summaries);
#else
    for (int k = 0; k < 20 * 40; k++)
    {
        threadIdx.x = k / 40;
        threadIdx.y = k % 40;
        Kernel(d_obs, d_bucket, d_summaries);
    }
#endif      

    cudaMemcpy(summary, d_summary, MAX_BUCKETS * sizeof(Info), cudaMemcpyDeviceToHost);

    time_t end;
    time(&end);
    cout << "Finished calculations in " << difftime(end, start) << "s" << endl;
    cin.get();
    return 0;
}

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

__global__ void Kernel(Observation * d_obs, 
                         int * d_bucket,
                         double * values,
                         Info ** d_summaries)
{
    Info * summary = d_summaries[blockIdx.x * 40 + blockIdx.y];

    __shared__ Info working_summary[1024];
    working_summary[threadIdx.x] = summary[threadIdx.x];
    __syncthreads();

    for (int i = 0; i < MAX_OBS; i++)
    {
        if (d_bucket[i] != threadIdx.x) continue;
        if (d_obs[i].core[blockIdx.x] < values[blockIdx.y])
            working_summary[threadIdx.x].left++;
        else
            working_summary[threadIdx.x].right++;
    }
    __syncthreads();

    summary[threadIdx.x] = working_summary[threadIdx.x];
} 

Это займет 18 секунд для <<<dim(20, 40, 1), 1000>>> и 172 для <<<dim(20,40,10), 1000>>> --- что хуже, чем один поток ЦП и линейно увеличивается в количестве параллельных задач.

Answer 1

Используемая вами плата K520 имеет два графических процессора, каждое из которых имеет 8 потоковых мультипроцессоров, с пиковой пропускной способностью ~160 ГБ / с на каждый графический процессор. С помощью приведенного выше кода вы должны быть ограничены этой пропускной способностью и должны рассчитывать на получение как минимум 100 ГБ / с на графический процессор (хотя я бы хотел запустить один графический процессор). Может быть, вы не сможете поразить его, может быть, вы его победите, но это хорошая цель, к которой нужно стремиться.

Количество блоков

Первое, что нужно сделать, это исправить параметры запуска. Эта строка:

Kernel<<<1, dim3(20, 40, 1)>>>(d_obs, d_bucket, d_summaries);

означает, что вы запускаете 1 блок CUDA из 800 потоков. Это далеко не достаточно параллелизма для графического процессора. Вам нужно как минимум столько же блоков, сколько потоковых мультипроцессоров (т. Е. 8), предпочтительно значительно больше (т. Е. 100+). Это даст вам значительное улучшение производительности. Параллелизм с 800 путями только недостаточно для графического процессора.

Разбросанные записи

Графические процессоры могут быть довольно чувствительными к шаблонам доступа. Следующий код:

summary[d_bucket[i]].left++;

делает разбросанную 4-байтовую запись в резюме. Транзакции с разрозненной памятью являются дорогостоящими на графическом процессоре, и для разумной работы с кодами, связанными с памятью, их следует избегать. Что мы можем сделать с этим в этом случае? Решение, на мой взгляд, добавить больше параллелизма. Вместо того, чтобы иметь сводку по каждому потоку, иметь сводку по каждому блоку. Каждый поток может работать на подмножестве диапазона 0...MAX_OBSи может увеличивать сводный массив на уровне блоков, который должен находиться в shared memory, В конце ядра вы можете записать результат обратно в глобальную память. К счастью, это также решает проблему отсутствия параллелизма, отмеченную выше!

Что дальше?

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