Копирование памяти двумя ядрами CUDA - почему скорость отличается?

Question

Может кто-нибудь помочь мне понять разницу в производительности между ядрами memCopy2dA и memCopy2dB?

Предполагается, что они копируют 2D-данные с размерами xLen, yLen из одного места в другое, но используют разные стратегии:

• , когда memCopy2dA использует блоки / потокиохватывать все 2D-пространство, так как предполагается, что это ядро ​​копирует только одну точку данных

• , когда используется memCopy2dB, блоки / потоки создаются только для одной всей строки X, а затем каждое ядро ​​проходит циклНаправление Y для копирования всех данных.

Согласно профилировщику (nvvp), в обоих случаях структура памяти доступа к GPU составляет 100%, а размер X достаточно большой, чтобы насыщать устройство для ядра "B" (Титан Х, 24см).К сожалению, ядро ​​«B» работает медленнее, и на моей машине результат:

GB/s: 270.715
GB/s: 224.405

Дополнительный вопрос: возможно ли даже приблизиться к теоретическому пределу пропускной способности памяти, который составляет 336,48 ГБ / с (3505 МГц * 384 бит *)2/8)По крайней мере, мои тесты показывают, что максимальное значение всегда составляет около 271-272 ГБ / с.

Код теста:

#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <iostream>
#include <chrono>

template<typename T>
__global__ void memCopy2dA(T *in, T *out, size_t xLen, size_t yLen) {
    int xi = blockIdx.x * blockDim.x + threadIdx.x;
    int yi = blockIdx.y * blockDim.y + threadIdx.y;
    if (xi < xLen && yi < yLen) {
        out[yi * xLen + xi] = in[yi * xLen + xi];
    }
}

template<typename T>
__global__ void memCopy2dB(T *in, T *out, size_t xLen, size_t yLen) {
    int xi = blockIdx.x * blockDim.x + threadIdx.x;
    if (xi < xLen) {
        size_t idx = xi;
        for (int y = 0; y < yLen; ++y) {
            out[idx] = in[idx];
            idx += xLen;
        }
    }
}

static void waitForCuda() {
    cudaDeviceSynchronize();
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) printf("Error: %s\n", cudaGetErrorString(err));
}

int main() {
    typedef float T;

    size_t xLen = 24 * 32 * 64; //49152
    size_t yLen = 1024;
    size_t dataSize = xLen * yLen * sizeof(T);

    T *dInput;
    cudaMalloc(&dInput, dataSize);
    T *dOutput;
    cudaMalloc(&dOutput, dataSize);

    const int numOfRepetitions = 100;
    double gigabyte = 1000 * 1000 * 1000;
    {
        dim3 threadsPerBlock(64, 1);
        dim3 numBlocks((xLen + threadsPerBlock.x - 1) / threadsPerBlock.x,
                       (yLen + threadsPerBlock.y - 1) / threadsPerBlock.y);

        auto startTime = std::chrono::high_resolution_clock::now();
        for (int i = 0; i < numOfRepetitions; ++i) {
            memCopy2dA <<< numBlocks, threadsPerBlock >>> (dInput, dOutput, xLen, yLen);
            waitForCuda();
        }
        auto stopTime = std::chrono::high_resolution_clock::now();
        std::chrono::duration<double> elapsed = stopTime - startTime;
        std::cout << "GB/s: " << (2 * dataSize * numOfRepetitions) / elapsed.count() / gigabyte << std::endl;
    }
    {
        dim3 threadsPerBlock(64);
        dim3 numBlocks((xLen + threadsPerBlock.x - 1) / threadsPerBlock.x);

        auto startTime = std::chrono::high_resolution_clock::now();
        for (int i = 0; i < numOfRepetitions; ++i) {
            memCopy2dB <<< numBlocks, threadsPerBlock >>> (dInput, dOutput, xLen, yLen);
            waitForCuda();
        }
        auto stopTime = std::chrono::high_resolution_clock::now();
        std::chrono::duration<double> elapsed = stopTime - startTime;
        std::cout << "GB/s: " << ((2 * dataSize * numOfRepetitions) / elapsed.count()) / gigabyte << std::endl;
    }

    cudaFree(dInput);
    cudaFree(dOutput);

    return 0;
}

, скомпилированный с:

nvcc -std=c++11 memTest.cu -o memTest

Answer 1

Я нашел решение, как ускорить работу ядра memCopy2dB.Вот тесты, выполненные на 1080Ti (TITAN X больше не доступен для меня).Код из части вопроса дает следующие результаты:

GB/s: 365.423
GB/s: 296.678

более или менее это та же процентная разница, что наблюдалась ранее на Titan X. И теперь модифицированное ядро ​​memCopy2dB выглядит так:

template<typename T>
__global__ void memCopy2dB(T *in, T *out, size_t xLen, size_t yLen) {
    int xi = blockIdx.x * blockDim.x + threadIdx.x;
    if (xi < xLen) {
        size_t idx = xi;
        for (int y = 0; y < yLen; ++y) {
            __syncthreads();  // <------ this line added
            out[idx] = in[idx];
            idx += xLen;
        }
    }
}

Существует много информации о том, насколько важны операции объединенной памяти на уровне деформации, когда все потоки в деформации должны получать доступ к одинаковым выровненным сегментам памяти.Но кажется, что синхронизация деформаций в блоке делает возможным объединение на уровне между деформациями, возможно, с использованием большей ширины шины памяти на разных графических процессорах <- это только мое «объяснение» этой проблемы, поскольку я не смог найти никакой литературы по этому вопросу. </p>

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

GB/s: 365.255
GB/s: 352.026

Так что даже если выполнение кода замедляетсяблагодаря синхронизации мы получаем гораздо лучшие результаты.Я попробовал эту технику на некоторых моих кодах, которые обрабатывали данные в виде шаблона доступа к memCopy2dB, и это дало мне хорошее ускорение.