Вычисления графического процессора OpenCL/CUDA со скоростью звука. Любой метод достаточно быстрый, чтобы считывать данные с графического процессора один раз для каждого аудиобуфера? (т.е. минимум ~43 FPS)

Ну, я не знаю о лучших практиках в отношении. задержка, особенно в OpenCL. Но я могу предложить простой тест для измерения времени прохождения туда и обратно.

На самом деле я делаю здесь только две вещи:

1. Используйте двойную буферизацию, чтобы графический процессор был загружен.
2. Используйте закрепленную память хоста и не делайте явных копий.

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

На моем оборудовании (ноутбук Nvidia T1200 и настольный компьютер RTX 3090 под управлением Linux) эта установка выполняет двустороннюю передачу 1024 выборок за 15 микросекунд последовательно после первых одного или двух вызовов ядра.

Вот код:

#include <cuda_runtime.h>

#include <algorithm>
// using std::fill_n
#include <cstdio>
// using std::printf
#include <chrono>
// using std::steady_clock


/**
 * Simple input = output kernel
 */
__global__ void kernel(unsigned* out, const unsigned* in, int n)
{
    const int idx = blockDim.x * blockIdx.x + threadIdx.x;
    if (idx < n)
        out[idx] = in[idx];
}
/**
 * Creates a time stamp in microseconds
 *
 * No defined zero-time. Only useful for measuring relative time intervals
 */
unsigned current_time_us()
{
    using us = std::chrono::microseconds;
    return static_cast<unsigned>(std::chrono::duration_cast<us>(
            std::chrono::steady_clock::now().time_since_epoch()).count());
}
/** Fills the buffer with the current time stamp */
void fill_current_time(unsigned* buf, int n)
{
    std::fill_n(buf, n, current_time_us());
}

int main()
{
    int samples = 1024, repetitions = 100;
    int blocksize = 128;
    int gridsize = (samples + blocksize - 1) / blocksize;
    cudaStream_t stream;
    if (cudaStreamCreate(&stream))
        return 1;
    /*
     * We use pinned host memory that is directly accessible by the device and
     * the host for input and output transfer.
     * Two input and two output buffers for double-buffering
     */
    unsigned* transfer_bufs;
    if (cudaHostAlloc(&transfer_bufs, 4 * samples * sizeof(unsigned), 0))
        return 2;
    unsigned* input_bufs = transfer_bufs;
    unsigned* output_bufs = transfer_bufs + 2 * samples;
    /*
     * We use events for quick notification when a kernel is done without
     * having to synchronize the stream
     */
    cudaEvent_t output_avail[2];
    for(cudaEvent_t& event: output_avail)
        if (cudaEventCreate(&event))
            return 3;
    /*
     * Initial fill of the first double buffer
     */
    fill_current_time(input_bufs, samples);
    kernel<<<blocksize, gridsize, 0, stream>>>(
            output_bufs, input_bufs, samples);
    if (cudaEventRecord(output_avail[0], stream))
        return 4;
    for(int i = 1; i < repetitions; ++i) {
        int cur_buf = i & 1;
        int last_buf = cur_buf ^ 1;
        int cur_offset = samples * cur_buf;
        int last_offset = samples * last_buf;
        /*
         * Schedule the next computation
         */
        fill_current_time(input_bufs + cur_offset, samples);
        kernel<<<blocksize, gridsize, 0, stream>>>(
                    output_bufs + cur_offset, input_bufs + cur_offset, samples);
        if (cudaEventRecord(output_avail[cur_buf], stream))
            return 5;
        /*
         * Wait for the previous computation
         */
        if (cudaEventSynchronize(output_avail[last_buf]))
            return 6;
        /*
         * Measure the time interval from filling the input buffer to
         * receiving it back in the output buffer
         */
        std::printf("RTT %u us\n", current_time_us() - output_bufs[last_offset]);
    }
    /*
     * Wait for the last computation. No need to check the results
     */
    if (cudaEventSynchronize(output_avail[(repetitions - 1) & 1]))
        return 7;
}

Выход:

RTT 94 us
RTT 22 us
RTT 12 us
RTT 15 us
RTT 15 us
RTT 15 us
RTT 15 us
RTT 15 us
...

Однако я должен также отметить, что использование всего графического процессора только с 1024 сэмплами кажется практически невозможным. Даже один мультипроцессор имеет больше потоков! Таким образом, хотя задержка передачи не является проблемой, фактическое использование вычислительных ресурсов без увеличения размера буфера будет проблемой.

Но я не знаю, возможно, вы смешиваете 32 источника входного сигнала по 1024 сэмпла каждый. Кстати, увеличение передачи в 32 раза в моих тестах увеличивает RTT только до 60 мкс.

Обычные копии DMA

Вот модифицированная версия, которая использует cudaMemcpyAsync вместо прямого доступа к закрепленной памяти хоста. RTT увеличивается незначительно — до 25 мкс на ноутбуке и 40–50 мкс на настольном компьютере. Поэтому я действительно не знаю, откуда берется цитируемая информация о механизме DMA, имеющем задержку в миллисекундах, и к какому оборудованию она применима.

#include <cuda_runtime.h>

#include <algorithm>
// using std::fill_n
#include <cstdio>
// using std::printf
#include <chrono>
// using std::steady_clock


/**
 * Simple input = output kernel
 */
__global__ void kernel(unsigned* out, const unsigned* in, int n)
{
    const int idx = blockDim.x * blockIdx.x + threadIdx.x;
    if (idx < n)
        out[idx] = in[idx];
}
/**
 * Creates a time stamp in microseconds
 *
 * No defined zero-time. Only useful for measuring relative time intervals
 */
unsigned current_time_us()
{
    using us = std::chrono::microseconds;
    return static_cast<unsigned>(std::chrono::duration_cast<us>(
            std::chrono::steady_clock::now().time_since_epoch()).count());
}
/** Fills the buffer with the current time stamp */
void fill_current_time(unsigned* buf, int n)
{
    std::fill_n(buf, n, current_time_us());
}

int main()
{
    int samples = 1024, repetitions = 100;
    int blocksize = 128;
    int gridsize = (samples + blocksize - 1) / blocksize;
    cudaStream_t in_stream, compute_stream, out_stream;
    for(cudaStream_t* stream: {&in_stream, &compute_stream, &out_stream})
        if (cudaStreamCreate(stream))
            return 1;
    /*
     * Pinned host memory for data transfer. Double buffering
     */
    unsigned* host_bufs;
    if (cudaHostAlloc(&host_bufs, 4 * samples * sizeof(unsigned), 0))
        return 2;
    unsigned* host_input_bufs = host_bufs;
    unsigned* host_output_bufs = host_bufs + 2 * samples;
    /*
     * Device-side memory. Again double-buffering
     */
    unsigned* dev_bufs;
    if (cudaMalloc(&dev_bufs, 4 * samples * sizeof(unsigned)))
        return 3;
    unsigned* dev_input_bufs = dev_bufs;
    unsigned* dev_output_bufs = dev_bufs + 2 * samples;
    /*
     * We use events for quick notification when a kernel is done without
     * having to synchronize the stream
     */
    cudaEvent_t events[6];
    for(cudaEvent_t& event: events)
        if (cudaEventCreate(&event))
            return 4;
    cudaEvent_t* in_avail = events;
    cudaEvent_t* out_avail = events + 2;
    cudaEvent_t* out_on_host = events + 4;
    auto compute = [=](int i) {
        int cur_buf = i & 1;
        int cur_offset = samples * cur_buf;
        fill_current_time(host_input_bufs + cur_offset, samples);
        if (cudaMemcpyAsync(dev_input_bufs + cur_offset,
                host_input_bufs + cur_offset, samples * sizeof(unsigned),
                cudaMemcpyDefault, in_stream))
            return 5;
        if (cudaEventRecord(in_avail[cur_buf], in_stream))
            return 6;
        if (cudaStreamWaitEvent(compute_stream, in_avail[cur_buf], 0))
            return 7;
        kernel<<<blocksize, gridsize, 0, compute_stream>>>(
                dev_output_bufs + cur_offset, dev_input_bufs + cur_offset,
                samples);
        if (cudaEventRecord(out_avail[cur_buf], compute_stream))
            return 8;
        if (cudaStreamWaitEvent(out_stream, out_avail[cur_buf], 0))
            return 9;
        if (cudaMemcpyAsync(host_output_bufs + cur_offset,
                dev_output_bufs + cur_offset, samples * sizeof(unsigned),
                cudaMemcpyDefault, out_stream))
            return 10;
        if (cudaEventRecord(out_on_host[cur_buf], out_stream))
            return 11;
        return 0;
    };
    /*
     * Initial fill of the first double buffer
     */
    if (compute(0))
        return 12;
    for(int i = 1; i < repetitions; ++i) {
        /*
         * Schedule next computation
         */
        if (compute(i))
            return 13;
        /*
         * Wait for previous
         */
        int last_buf = (i - 1) & 1;
        int last_offset = samples * last_buf;
        if (cudaEventSynchronize(out_on_host[last_buf]))
            return 14;
        /*
         * Measure the time interval from filling the input buffer to
         * receiving it back in the output buffer
         */
        std::printf("RTT %u us\n",
                current_time_us() - host_output_bufs[last_offset]);
    }
    /*
     * Wait for the last computation. No need to check the results
     */
    if (cudaEventSynchronize(out_on_host[(repetitions - 1) & 1]))
        return 15;
}