Почему производительность cuFFT страдает с перекрывающимися входами?

Question

Я экспериментирую с использованием функции обратного вызова cuFFT для преобразования входного формата «на лету» (например, вычисления БПФ 8-битных целочисленных входных данных без предварительного преобразования входного буфера в float). Во многих моих приложениях мне нужно вычислить с перекрытием FFT на входном буфере, as described in this previous SO question. Как правило, смежные БПФ могут перекрываться на 1/4 до 1/8 длины БПФ.

cuFFT, с его интерфейсом FFTW, явно поддерживает это via the idist parameter of the cufftPlanMany() function. В частности, если я хочу рассчитать БПФ размером 32768 с перекрытием 4096 выборок между последовательными входами, я бы установил idist = 32768 - 4096. Этот делает правильной работой в том смысле, что он дает правильный выход.

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

1. Явного сказать CUFFT о перекрывающей природе ввода: установить idist = nfft - overlap как я описано выше. Установите функцию обратного вызова нагрузки, которая только делает преобразование от int8_t до float по мере необходимости на индекс буфера, предоставленный для обратного вызова.

2. Не говорите cuFFT о перекрывающемся характере ввода; ложь ему dset idist = nfft. Затем пусть функция обратного вызова обрабатывает перекрытие путем вычисления правильного индекса, который должен быть прочитан для каждого входа FFT.

A test program implementing both of these approaches with timing and equivalence tests is available in this GitHub gist. Я не воспроизводил все это здесь для краткости. Программа вычисляет партию из 1024 32768-точечных БПФ, которые перекрываются 4096 выборками; тип входных данных - 8-битные целые числа. Когда я запускаю его на моей машине (с Geforce GTX 660 GPU, используя CUDA 8.0 RC на Ubuntu 16.04), я получаю следующий результат:

executing method 1...done in 32.523 msec 
executing method 2...done in 26.3281 msec 

Метод 2 заметно быстрее, что я бы не ожидал. Посмотрите на реализаций функций обратного вызова:

Метод 1:

template <typename T> 
__device__ cufftReal convert_callback(void * inbuf, size_t fft_index, 
    void *, void *) 
{ 
    return (cufftReal)(((const T *) inbuf)[fft_index]); 
} 

Способ 2:

template <typename T> 
__device__ cufftReal convert_and_overlap_callback(void *inbuf, 
    size_t fft_index, void *, void *) 
{ 
    // fft_index is the index of the sample that we need, not taking 
    // the overlap into account. Convert it to the appropriate sample 
    // index, considering the overlap structure. First, grab the FFT 
    // parameters from constant memory. 
    int nfft = overlap_params.nfft; 
    int overlap = overlap_params.overlap; 
    // Calculate which FFT in the batch that we're reading data for. This 
    // tells us how much overlap we need to account for. Just use integer 
    // arithmetic here for speed, knowing that this would cause a problem 
    // if we did a batch larger than 2Gsamples long. 
    int fft_index_int = fft_index; 
    int fft_batch_index = fft_index_int/nfft; 
    // For each transform past the first one, we need to slide "overlap" 
    // samples back in the input buffer when fetching the sample. 
    fft_index_int -= fft_batch_index * overlap; 
    // Cast the input pointer to the appropriate type and convert to a float. 
    return (cufftReal) (((const T *) inbuf)[fft_index_int]); 
} 

Метод 2 имеет значительно более сложную функцию обратного вызова, один что даже включает в себя целочисленное деление по не компилируемому значению времени! Я ожидал бы, что это будет намного медленнее, чем метод 1, но я вижу обратное. Есть ли хорошее объяснение этому? Возможно ли, что cuFFT структурирует свою обработку по-разному, когда вход перекрывается, что приводит к ухудшению производительности?

Похоже, я должен быть в состоянии достичь производительности, которая совсем немного быстрее, чем метод 2, если расчеты индекса могут быть удалены из обратного вызова (но это потребует перекрытия, который будет указан в CUFFT).

Edit: После запуска моей тестовой программы под nvvp, я могу видеть, что CUFFT определенно, кажется, структурирования ее вычисления по-разному. Это трудно понять смысл имен символов ядра, но вызовы ядра ломаются, как это:

Метод 1:

1. __nv_static_73__60_tmpxft_00006cdb_00000000_15_spRealComplex_compute_60_cpp1_ii_1f28721c__ZN13spRealComplex14packR2C_kernelIjfEEvNS_19spRealComplexR2C_stIT_T0_EE: 3,72 мс
2. spRadix0128C::kernel1Tex<unsigned int, float, fftDirection_t=-1, unsigned int=16, unsigned int=4, CONSTANT, ALL, WRITEBACK>: 7,71 мс
3. spRadix0128C::kernel1Tex<unsigned int, float, fftDirection_t=-1, unsigned int=16, unsigned int=4, CONSTANT, ALL, WRITEBACK> : 12,75 мс (да, он активируется дважды)
4. __nv_static_73__60_tmpxft_00006cdb_00000000_15_spRealComplex_compute_60_cpp1_ii_1f28721c__ZN13spRealComplex24postprocessC2C_kernelTexIjfL9fftAxii_t1EEEvP7ComplexIT0_EjT_15coordDivisors_tIS6_E7coord_tIS6_ESA_S6_S3_: 7.49 мсек

Метод 2:

1. spRadix0128C::kernel1MemCallback<unsigned int, float, fftDirection_t=-1, unsigned int=16, unsigned int=4, L1, ALL, WRITEBACK>: 5,15 мс
2. spRadix0128C::kernel1Tex<unsigned int, float, fftDirection_t=-1, unsigned int=16, unsigned int=4, CONSTANT, ALL, WRITEBACK>: 12,88 мс
3. __nv_static_73__60_tmpxft_00006cdb_00000000_15_spRealComplex_compute_60_cpp1_ii_1f28721c__ZN13spRealComplex24postprocessC2C_kernelTexIjfL9fftAxii_t1EEEvP7ComplexIT0_EjT_15coordDivisors_tIS6_E7coord_tIS6_ESA_S6_S3_: 7,51 мс

Интересно, что это выглядит как CUFFT вызывает два ядра на самом деле вычислить БПФ с использованием метода 1 (когда cuFFT знает о перекрытии), но с методом 2 (где он не знает t хх, FFTs перекрываются), он выполняет задание только с одним. Для ядер, которые используются в обоих случаях, они, похоже, используют одни и те же параметры сетки между методами 1 и 2.

Не понимаю, почему здесь нужно использовать другую реализацию, тем более, что входной шаг istride == 1. Он должен просто использовать другой базовый адрес при извлечении данных на входе преобразования; я думаю, остальная часть алгоритма должна быть точно такой же.

Редактировать 2: Я вижу еще более странное поведение. Я случайно осознал, что если я не смогу самостоятельно уничтожить ручки cuFFT, я вижу различия в измеренной производительности. Например, я модифицировал тестовую программу, чтобы пропустить уничтожение дескрипторов cuFFT, а затем выполнить тесты в другой последовательности: метод 1, метод 2, затем метод 2 и метод 1 снова. Я получил следующие результаты:

executing method 1...done in 31.5662 msec 
executing method 2...done in 17.6484 msec 
executing method 2...done in 17.7506 msec 
executing method 1...done in 20.2447 msec 

Так производительность, кажется, изменяется в зависимости от того, есть ли другие планы CUFFT в существовании при создании плана для теста! Используя профилировщик, я вижу, что структура запуска ядра не изменяется между двумя случаями; ядра просто все работают быстрее. У меня нет разумного объяснения этого эффекта.

Answer 1

По предложению @llukas я опубликовал отчет об ошибке с NVIDIA по поводу проблемы (https://partners.nvidia.com/bug/viewbug/1821802, если вы зарегистрированы как разработчик). Они признали худшую работу с перекрывающимися планами. Фактически они указали, что конфигурация ядра, используемая в обоих случаях, является неоптимальной, и в конечном итоге они планируют улучшить ее. Нет ETA, но, скорее всего, не будет быть в следующем выпуске (8,0 был выпущен только на прошлой неделе). Наконец, они сказали, что с CUDA 8.0 не существует обходного пути для использования cuFFT более эффективного метода со строгими входами.

Answer 2

Если вы укажете нестандартные шаги (не имеет значения, если пакет/преобразование) cuFFT использует другой путь внутри.

ad edit 2: Это, скорее всего, GPU Boost, настраивающие часы на GPU.План CUFFT не влияют друг на друга

Пути, чтобы получить более стабильные результаты:

1. запуска прогрева ядра (ничего, что бы сделать полный работу GPU хорошо), а затем вашу проблему
2. увеличение размера партии
3. тест запустить несколько раз и взять среднее
4. блокировки часов ГПУ (на самом деле не возможно на GeForce - Tesla может это сделать)