CC BY
130
УДК 004.42'274:621.391.083.92 ББК З811.3-043:В161.911
НА. ГАЛАНИНА, А.Г. АЛЕКСЕЕВ, А.В. СЕРЕБРЯННИКОВ
ВЫЧИСЛЕНИЕ БЫСТРОГО ПРЕОБРАЗОВАНИЯ ФУРЬЕ С ИСПОЛЬЗОВАНИЕМ ТЕХНОЛОГИИ CUDA*
Ключевые слова: цифровая обработка сигналов, дискретное преобразование Фурье, быстрое преобразование Фурье, гетерогенное параллельное программирование, технология CUDA, архитектура Kepler, быстродействие.
В основе различных технологий исследования цифровых сигналов лежит дискретное преобразование Фурье, по возможности вычисляемое быстрыми методами. В статье предлагается решение задачи повышения скорости спектрального анализа сигналов способами гетерогенного параллельного программирования с использованием технологии Cuda. Определяется зависимость времени выполнения преобразования Фурье от размера задачи. Проводится анализ путей повышения эффективности вычислений. В качестве тестовых примеров рассматривается вычисление прямого преобразования Фурье на последовательностях, являющихся степенями двойки. Вычисление быстрого преобразования Фурье проводится с использованием библиотеки cuFFT. Для исследования используются графические ускорители GeForce GT 750M и GeForce GTX 670, изготовленные на базе архитектуры Kepler.
Быстрое преобразование Фурье (БПФ) является одним из наиболее востребованных алгоритмов цифровой обработки сигналов (ЦОС) и моделирования физических процессов [1, 2]. Большинство существующих процессоров ЦОС и устройств на ПЛИС эффективно работают только с выборками до 214 отсчетов. Обработка выборок большего размера, требуемая в ряде задач для получения высокой точности, традиционно выполнялась только на процессорах общего назначения. Благодаря развитию гетерогенного параллельного программирования появились возможности реализации алгоритма спектрального анализа с использованием графических адаптеров для работы с большими данными.
В настоящее время для гетерогенных вычислений широко используется технология CUDA (Compute Unified Device Architecture) от компании NVIDIA. Исследовательский интерес представляют реализация вычисления БПФ с ее использованием, анализ путей повышения эффективности вычислений и дальнейшая выработка практических рекомендаций.
Технология CUDA (архитектура Kepler). Гетерогенными называются вычисления, которые выполняются c использованием различных видов вычислительных устройств. В качестве таких устройств могут выступать процессоры общего назначения (GPP), процессоры цифровой обработки сигналов (DSP), графические процессоры (GPU), процессоры специального назначения и т.д. Ключевым аспектом является то, что эти устройства построены на архитектурах с различающимися наборами команд. Технология CUDA разработана компанией NVIDIA в 2006 г. и поддерживается несколькими поколениями графических ускорителей. В рамках исследования использовались два графических ускорите-
* Исследование выполнено при финансовой поддержке РФФИ и Кабинета Министров Чувашской Республики в рамках научного проекта № 17-47-210790 р_а.
ля - GeForce GT 750M и GeForce GTX 670, изготовленные на базе архитектуры Kepler (рис. 1).
На рис. 1 видно, что в состав потокового мультипроцессора (Streaming Multiprocessor, SMX), являющегося базовым элементом технологии CUDA и выполняющего основную работу, входят 192 ядра CUDA (Core). Помимо них имеется 64 блока для работы с вещественными числами двойной точности (Double Precision Unit, DP Unit) и 32 блока специального назначения (Special Function Unit, SFU). Кроме того, имеется 32 блока для доступа к видеопамяти (Load/Store, LDST).
CUspntrh Blspalch * M
Warp Scheduler warp sehedul и Warp scftotiuler
и s ре le h [îlspalch Dispatch Dl&patchi Dispatch Dfeparte!)
Register File (65.536 x 32-bill
i * Cwu tin» * * * Cm Core + IftST - + » + Co* Cor* Con + * 0PUfÜ Con Con ^ 4 Can * IDfST 1Г- SFIf
CWH Cor» Cur- DPUrtt Ce» Con C«*» jy.ÎMt U>TT SFU Coi* Core Core IQfU« Cora Cwe CK, 1Ю1Т 8FU
0« Сея* ■ ppyriil Co-с См* Cart РРЦЦ* Lft-вт sru С сиг Çore COT« ■ ■■ Core Ш* SFU
Corn Сен* Con И* ÜWt СО'Ч Cora Core- ИРЦяИ LOVÖT EFU Сем* Coro Cor® и— - 9.ШВ1 SFU
И Cor* с*. 0PIMH; Соч ■ ■■ iQvST sru Cor* Coro Com QPlMt Core Cere Corn ithtT SFU
cufft сея* ce» OPUMt Сош ■ ■■ IKST SFU CO>* Cor« CO« IGNÜ Urtft Cm Co« COr« ItHST SFU
сек* CS» WIM с спи C« .Щ L&flT SFU Cor* Cor« COT« QPUtrit Согв C«» Car« LD1ST SFU
Cora Cor* Cüru DP Lfm il Сеч Cora Core ÇHMNéM, IWT EFU Cor* Cor* Com lOpiïMI Core Ce« Corn IMT SFU
Cora Cor* Си* WltoH. Сот® Cora Cor» IDPUNM IIXST SFU Ce*» Cora Cora |fiMdmi| Cora С04* Corn (Mr SFU
Cora со* C-Ure --И 1И1 SFU Core Cor« Core Core Core COT» LOT SFW
Gore C<*» QFUW Cp.» ■■H 1ИТ EFU Core Core Core Ityilnl Core Core - IWST SFU
COfd ce™ COM WUMt Card cor« Ca«' LKBt 5FU cm cor* cor« ИНН ШШ Cor* M 1Ш5Г SFU
Cora cars Coro WUintt Core cor« Cm МИШ IDfflT 5FU Co«* Cora Core ÉfrUnlt Cors car* Com ItHST SFU
Coro bi Cor* Core Core- DP'UMt ЮТГ SFU Con- Cora Core âPIMt Core Cor* Core ЮТ sru
Corn Core Coro DPUMt Core Con Cere- DPUfiK UWT SFU Core Cons Core Core Core Care SFU
С» Cef* № Urn См M с» ¡Д IMT SFU Cor* Cor* Core 0*U* Core Cor* Cor« imsr SFU
Interconnect Network M KB Shared Memory IH Cael* 49 № Read-Only Data Cache
Рис. 1. Архитектура SMX (Kepler) [7]
Все эти блоки используются для параллельных SIMD-вычислений (Single Instruction Multiple Data) по 32 потока в каждой группе. Такая группа потоков получила название warp. В рассматриваемой архитектуре одновременно могут выполняться до 4 warp (диспетчеризацией групп занимаются специальные планировщики Warp Scheduler). Каждый планировщик может обрабатывать до 16 warp, но активен в каждый момент времени может быть только один.
Мультипроцессор содержит:
- 65 536 32-разрядных регистров, используемых для хранения значений переменных потоков;
- 64 Кб разделяемой памяти, через которую потоки могут обмениваться данными между собой;
- 48 Кб памяти, доступной только для чтения (используется для хранения констант).
Также существует возможность работы с текстурной памятью (Texture, Tex). На одном графическом ускорителе, как правило, устанавливается несколько таких потоковых мультипроцессоров. Кроме того, графические ускорители отличаются количеством и типом видеопамяти.
Использование библиотеки cuFFT. Библиотека cuFFT с широким спектром возможностей от компании NVIDIA [5] основана на алгоритме БПФ Кули - Тьюки, использующем принцип «разделяй и властвуй», и предназначена для эффективного вычисления БПФ. Будем применять ее для прямого преобразования на последовательностях, являющихся степенями двойки.
Программирование на CUDA C. Для программирования в технологии CUDA используется синтаксис языка С. Для разработки программ необходим программный пакет CUDA Toolkit, а для их выполнения - драйвер устройства с поддержкой CUDA (обычно устанавливается в комплекте с CUDA Toolkit). В состав CUDA Toolkit входит прикладной программный интерфейс среды выполнения - CUDA Runtime API, который представляет собой набор С/С++ функций для выполнения на основной ЭВМ, предназначенных для получения информации об устройстве, управления его конфигурацией, запуска вычислительных задач и т.п.
В состав программного пакета входит специальный компилятор языка С - NVCC. Файлы, предназначенные для компилирования при помощи NVCC, обычно имеют расширение «.cu».
В состав программы входят участки кода (функции), которые выполняются на основной ЭВМ, и участки кода, которые выполняются на устройстве (видеоадаптере).
Текст программы на CUDA C:
#include <cuda.h> #include <cuda runtime.h> #include <cufft.h>
#include "device launch parameters.h" #include <stdio.h> #include <vector> #include <iostream> #include <fstream>
const float PI_F = 3.14159265358979f;
// Макрос для проверки на ошибку, при вызове функций CUDA #define CUDA_CHECK(call) \ if((call) != cudaSuccess) { \
cudaError t err = cudaGetLastError(); \
printf( "CUDA error calling \"%s\", error code is %d\n", #call,
err); \
printf( "Error %d: %s\n", err, cudaGetErrorString(err)); \ exit(-1); }
int main(int argc, char* argv[]) {
// определяем размер последовательности из аргументов if (argc != 2) return -1; size_t N = 1 << atol(argv[1]); printf("Sequence length: %d\n", N);
// генерируем входной сигнал и сохраняем его в файл "source.dat" std::ofstream myfile; myfile.open("source.dat");
cufftReal *source signal = (cufftReal*)malloc(N * sizeof(cufftReal)); float w = 2 * PI_F / N;
for (size_t i = 0; i < N; ++i) {
source signal[i] = // синусоида с амплитудой A1 и частотой f1
0.3 * sinf(N / 17 * w * static_cast<float>(i)) + // синусоида с амплитудой A2 и частотой f2
0.3 * sinf(N / 3 * w * static_cast<float>(i)) + // синусоида с амплитудой A2 и частотой f2
0.5 * sinf(N / 5 * w * static_cast<float>(i)) + // случайный шум
0.2 * rand() / static_cast<float>(RAND_MAX) - 0.1;
myfile << source signal[i] << "\n";
}
myfile.close();
// проверка на наличие устройств с поддержкой CUDA int deviceCount = 0;
CUDA_CHECK(cudaGetDeviceCount(&deviceCount));
if (deviceCount == 0) {
printf("There are no available device(s) that support CUDA\n"); return 0;
}
// используем первое устройство cudaSetDevice(0);
cudaEvent t evt1, evt2, evt3, evt4, evt5, evt6;
cudaEventCreate(&evt1);
cudaEventCreate(&evt2);
cudaEventCreate(&evt3);
cudaEventCreate(&evt4);
cudaEventCreate(&evt5);
cudaEventCreate(&evt6);
// выделение видеопамяти cudaEventRecord(evt1, 0); cufftComplex *device data; CUDA_CHECK(cuda^alloc((void**)&device_data, N * sizeof(cufftComplex)));
// копирование исходных данных в видеопамять cudaEventRecord(evt2, 0);
size t nBytes = N / 2 * sizeof(cufftReal); CUDA CHECK(cuda^emcpy(device data, source signal, nBytes, cudaMemcpyHostToDevice));
// обработка сигнала cudaEventRecord(evt3, 0); cufftHandle plan;
CUDA_CHECK(cufftPlan1d(&plan, N, CUFFT_R2C, 1)); CUDA_CHECK(cufftExecR2C(plan, (cufftReal*)device_data, device_data)); CUDA_CHECK(cufftDestroy(plan));
// копирование результата из видеопамяти cudaEventRecord(evt4, 0);
std::vector<cufftComplex> result(N / 2 + 1); CUDA_CHECK(cudaMemcpy(result.data(), device_data, sizeof(cufftComplex)*(N / 2 + 1), cudaMemcpyDeviceToHost)); // освобождаем видеопамять cudaEventRecord(evt5, 0); CUDA_CHECK(cudaFree(device_data));
// окончание работы cudaEventRecord(evt6, 0); cudaEventSynchronize(evt6);
// сохраняем результаты работы в файл "result.dat" myfile.open("result.dat"); for (size_t i = 0; i < (N / 2 + 1); ++i) myfile << sqrt(result[i].x * result[i].x + result[i].y * result[i].y) / N << "\n";
myfile.close();
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, evt1, evt2); printf("Video memory allocation: %g ms\n", elapsedTime); cudaEventElapsedTime(&elapsedTime, evt2, evt3);
printf("Copying source data to video memory: %g ms\n", elapsedTime); cudaEventElapsedTime(&elapsedTime, evt3, evt4); printf("Calculation: %g ms\n", elapsedTime); cudaEventElapsedTime(&elapsedTime, evt4, evt5);
printf("Copying results from video memory: %g ms\n", elapsedTime);
cudaEventElapsedTime(&elapsedTime, evt5, evt6);
printf("Video memory free: %g ms\n", elapsedTime);
cudaEventElapsedTime(&elapsedTime, evt1, evt6);
printf("Total time: %g ms\n", elapsedTime);
cudaEventDestroy(evt1);
cudaEventDestroy(evt2);
cudaEventDestroy(evt3);
cudaEventDestroy(evt4);
cudaEventDestroy(evt5);
cudaEventDestroy(evt6);
return 0;
На рис. 2. представлены официальные тесты производительности библиотеки от компании NVIDIA для одинарной и двойной точности вычислений [4].
Экспериментальная часть. В качестве входных данных будем использовать смесь трех синусоидальных сигналов с разными амплитудами и шума. Проверку выходных результатов будем осуществлять по спектру сигнала, на котором теоретически ожидаются три пика, соответствующие частотам первоначальных синусоид. Для построения спектрограммы воспользуемся программой gnuplot 5.0 [6].
Приведем характеристики используемых устройств:
Device: "GeForce GT 750M"
CUDA Driver Version / Runtime Version 8.0 / 8.0
CUDA Capability Major/Minor version number: 3.0
Total amount of global memory: 4096 MBytes (4294967296 bytes)
(2) Multiprocessors, (192) CUDA Cores/MP: 384 CUDA Cores
GPU Max Clock rate: 1085 MHz (1.09 GHz)
Memory Clock rate: 900 Mhz
Memory Bus Width: 128-bit
CUDA Device Driver Mode (TCC or WDDM): WDDM (Windows Display Driver Model)
Device: "GeForce GTX 670"
CUDA Driver Version / Runtime Version 8.0 / 8.0
CUDA Capability Major/Minor version number: 3.0
Total amount of global memory: 2048 MBytes (2147483648 bytes)
(7) Multiprocessors, (192) CUDA Cores/MP: 1344 CUDA Cores
GPU Max Clock rate: 1059 MHz (1.06 GHz)
Memory Clock rate: 3004 Mhz
Memory Bus Width: 256-bit
CUDA Device Driver Mode (TCC or WDDM): WDDM (Windows Display Driver Model)
cuFFT: up to 700 GFLOPS
1D Complex, BatchedFFTs Used in Audio Processing and as a Foundation for 2D and 3D FFTs
Single Precision
Double Precision
7 9 11 13 16 17 19 21 23 25 Iog2[transform_size)
250 200
I
i 150
i
100 50
7 9 11 13 IS 17 19 21 23 25 log2(transform_size)
Performance may vary based on Diversion and motherboard confier -лМпп
« cuFFT 6.0 en K4lk,ECC ON. Î2M elements, input and output data on de^ee
cuFFT: Consistently High Performance
800 700 600 и 500
CL
3 400 W 300 200 100 о
Single Precision
100 10.000 1,000.000 100,000,000 Transform Size
300 , 25« -200 -
tn a. 3 160 IL
et
100 50 0
Double Precision
100 10,000 1,000.000 100,000,000 Transform Size
-Powers Ы2 —^Powers of 3 Powers of 5 -Powers oF7
mance may vary based on Oi version and motherboard -тГи;, чМлп ■ r uFFT 6.0 nnK40c,ECC ON, 2ВМ-13ГЛ elements input and output data on desire
Рис. 2. Официальные тесты производительности библиотеки cuFFT [5]
Отметим, что используемые устройства различаются объемом памяти, типом шины памяти, а также количеством мультипроцессоров.
В табл. 1 и 2 приведены результаты испытаний: определены зависимости времени выполнения преобразования Фурье от размера входных данных для обоих устройств. Заметим, что продолжительность операции «выделение видеопамяти» не зависит от размерности данных, а операции копирования и
освобождения памяти зависят от размера данных линейно. Интересно отметить, что время самого преобразования нелинейно зависит от размера данных. Это можно объяснить тем, что чем больше размер задачи, тем больше коэффициент использования ядер. Также можно отметить, что большее количество ядер дает выигрыш только на задачах большего размера (см. рис. 3).
Таблица 1
Зависимость времени вычисления от размера последовательности (GeForce GT 750M), мc
Показатели Размерность преобразования (степень 2)
20 21 22 23 24 25 26 27
Выделение
видеопамяти 0,001024 0,001024 0,001024 0,001024 0,000992 0,00112 0,001024 0,000992
Копирование
исходных дан-
ных в видеопа-
мять 2,20333 2,43507 3,6991 5,84538 9,50858 17,6326 33,8018 64,0628
Вычисление
преобразования 250,174 255,708 281,294 278,413 376,489 375,963 514,677 837,595
Копирование
результатов
из видеопамяти 1,10218 2,1575 4,81293 8,44077 14,7201 29,6889 61,3887 130,991
Освобождение
видеопамяти 0,520288 0,794912 1,10707 1,77584 3,00832 6,00768 10,6497 21,9966
Общее время 254,001 261,097 290,914 294,476 403,727 429,293 620,518 1054,65
Таблица 2
Зависимость времени вычисления от размера последовательности (GeForce GTX 670), мc
Показатели Размерность преобразования (степень 2)
20 21 22 23 24 25 26 27
Выделение
видеопамяти 0,001088 0,00112 0,001152 0,001088 0,00112 0,001088 0,00112 0,001088
Копирование
исходных дан-
ных в видеопа-
мять 2,62835 3,71878 4,42819 6,87091 11,518 23,3011 61,5019 119,947
Вычисление
преобразования 277,942 276,493 284,707 285,379 303,753 337,943 402,187 612,738
Копирование
результатов
из видеопамяти 1,36029 1,90547 5,19194 7,62822 18,7774 40,571 55,102 109,396
Освобождение
видеопамяти 0,89744 0,686432 1,37133 1,6361 3,33808 6,70675 14,4116 27,4431
Общее время 282,829 282,805 295,7 301,515 337,387 408,523 533,203 869,525
1100 1000 900
u 800 * 700
I 600
* 500 400 300 200
20 21 22 23 24 25 26 27
Размерность преобразования (степень 2)
GeForce GT 750М — — — GeForce GTX 670 Рис. 3. Зависимость времени вычисления от размера задачи
На рис. 4 представлен спектр исследуемого сигнала, полученный на основе БПФ. Как и ожидалось, на нем присутствуют три пика, соответствующие частотам синусоид во входном сигнале.
0.14
0.12 0.1
го сС
£ 0.08
s ^
3 0.06 0.04 0.02 0
0 100000 200000 300000 400000 500000 600000
Номер отсчета Рис. 4. Спектр выходного сигнала на основе
БПФ (220 отсчетов)
Выводы. Библиотека cuFFT от компании NVIDIA предоставляет широкие возможности для вычисления БПФ. Данная библиотека является оптимизированной под вычисления с использованием графических ускорителей с поддержкой технологии CUDA. Она очень проста в использовании и требует минимальных усилий со стороны разработчика при работе с ней.
Зависимость времени выполнения преобразования Фурье от размера задачи носит нелинейный характер, что объясняется неравномерной загрузкой вычислительных ядер CUDA. В пользу этой гипотезы говорит и тот факт, что наличие большего числа ядер не дает заметного увеличения на задачах малой размерности.
"result.dat" i i
J_I__I__I_L
Визуализация результатов БПФ в виде спектра выходного сигнала вызывает сложности при размерностях большой длины. Используемая программа Gnuplot «справляется» только с данными меньше одного миллиона записей. Другие программы для визуализации результатов сталкиваются с подобными трудностями.
Таким образом, для достижения максимальной производительности необходимо выполнить три условия: максимально использовать доступные ядра, эффективно работать с памятью и обеспечить эффективное выполнение инструкций.
Литература
1. Алексеев А.Г., Йовенко А.Р. Параллельное программирование. Чебоксары: Изд-во Чуваш. ун-та, 2015. 194 с.
2. Галанина Н.А. Синтез функциональных модулей БПФ // Вестник Чувашского университета. 2005. № 2 . С. 124-127.
3. Галанина Н.А., Песошин В.А., Иванова Н.Н. Разработка устройств цифровой фильтрации и спектрального анализа с индексированием данных в системе остаточных классов // Вестник Чувашского университета. 2014. № 2. С. 93-97.
4. CuFFT [Электронный ресурс]. URL: https://developer.nvidia.com/cufft (дата обращения 22.12.2016).
5. CuFFT User Guide [Электронный ресурс]. URL: http://docs.nvidia.com/cuda/cufft/in-dex.html (дата обращения 22.12.2016).
6. Gnuplot tutorial [Электронный ресурс]. URL: http://www.gnuplot.info/docs_5.0/gnup-lot.pdf (дата обращения 22.12.2016).
7. Kepler Compute Architecture White Paper [Электронный ресурс]. URL: http://in-ternational.download.nvidia.com/pdf/kepler/NVIDIA-Kepler-GK110-GK210-Architecture-Whitepaper.pdf (дата обращения 22.12.2016).
ГАЛАНИНА НАТАЛИЯ АНДРЕЕВНА - доктор технических наук, профессор кафедры математического и аппаратного обеспечения информационных систем, Чувашский государственный университет, Россия, Чебоксары ([email protected]).
АЛЕКСЕЕВ АЛЕКСАНДР ГЕОРГИЕВИЧ - старший преподаватель кафедры вычислительной техники, Чувашский государственный университет, Россия, Чебоксары ([email protected]).
СЕРЕБРЯННИКОВ АЛЕКСАНДР ВЛАДИМИРОВИЧ - кандидат технических наук, доцент кафедры промышленной электроники, Чувашский государственный университет, Россия, Чебоксары ([email protected]).
N. GALANINA, A. ALEKSEEV, A. SEREBRYANNIKOV
CALCULATION OF FOURIER TRANSFORM USING CUDA TECHNOLOGY
Key words: digital signal processing (DSP), fast Fourier transform (FFT), heterogeneous parallel programming, CUDA technology, Kepler architecture, speed.
The discrete Fourier transform, possibly computed by fast methods, underlies the various technologies for studying digital signals. The paper proposes the solution of the problem of increasing the speed of spectral analysis of signals by means of heterogeneous parallel programming using Cuda technology. The dependence of the execution time of the Fourier transform on the size of the problem is determined. As test examples, the calculation of the direct Fourier transform on sequences that are powers of two is chosen. The FFT was calculated using the cuFFT library. For the study, GeForce GT 750M and GeForce GTX 670 graphics accelerators were used, both based on the Kepler architecture.
References
1. Alekseev A.G., Iovenko A.R. Parallel'noe programmirovanie [Parallel Programming]. Cheboksary, Chuvash University Publ., 2015, 194 p.
2. Galanina N.A. Sinthes tablichnykh moduley BPF v SOK [Synthesis of tabular FFT modules in RNS]. Vestnik Chuvashskogo universiteta, 2005, no. 2, pp. 124-127.
3. Galanina N.A., Pesoshin V.A., Ivanova N.N. Razrabotka ustroistv tsifrovoi fil'tratsii i spektral'nogo analiza s indeksirovaniem dannykh v sisteme ostatochnykh klassov [Development of devices of Digital Filtration and Spectral Analysis with the index of the data in Residue Number System]. Vestnik Chuvashskogo universiteta, 2014, no. 2, pp. 93-97.
4. CuFFT. Available at: https://developer.nvidia.com/cufft (Access Date 22 Dec. 2016).
5. CuFFT User Guide. Available at: http://docs.nvidia.com/cuda/cufft/index.html (Access Date 22 Dec. 2016).
6. Gnuplot tutorial. Available at: http://www.gnuplot.info/docs_5.0/gnuplot.pdf (Access Date 22 Dec. 2016).
7. Kepler Compute Architecture White Paper. Available at: http://international.download.nvi-dia.com/pdf/kepler/NVIDIA-Kepler-GK110-GK210-Architecture-Whitepaper.pdf (Access Date 22 Dec. 2016).
GALANINA NATALIYA - Doctor of Technical Sciences, Professor, Information Systems Math and Hardware Department, Chuvash State University, Russia, Cheboksary ([email protected]).
ALEKSEEV ALEKSANDER - Senior Lecturer, Department of Computer Science, Chuvash State University, Russia, Cheboksary ([email protected]).
SEREBRYANNIKOV ALEXANDER - Candidate of Technical Sciences, Associate Professor of Industrial Electronics Department, Chuvash State University, Russia, Cheboksary (alex-silver@mail. ru).
Ссылка на статью: Галанина Н.А., Алексеев А.Г., Серебрянников А.В. Вычисление быстрого преобразования Фурье с использованием технологии CUDA // Вестник Чувашского университета. - 2018. - № 1. - С. 88-97.
