Tento článek vznikl na základě samostatného projektu pro CVUT v Praze. Náplní bylo vytvořit FFT Algoritmus běžící na paralelní platformě (GPU) v jazyce OpenCL. Nicméně výsledkem nakonec byla verze pro jazyky C/C++, CUDA a OpenCL. Ideálním nasazením algoritmu by bylo vyhledávání GPS signálu. Nicméně Algoritmus jako takový, přestože je rychlejší než CPU verze (Jak si v závěru ukážeme) je stále příliš pomalý pro tento výpočet.

  • This article was created on the basis of my semestral work at CVUT in Prague. The goal was to implement Fast Fourier Trasform Algorithm on a parallel platform in OpenCL programming language. But as a result, I have implemented the algorithm both in OpenCL and CUDA. Though it was successfully implemented, it was still too slow for its purpose: To use it as a part of other GPS signal searching algorithms. (It should have been designed to run on mobile platform instead of high-end  PS GPU) …

 

Jaguar_longview

Úvodem | Intro

Algoritmus převádí ovzorkovaný vstupní signál na diskrétní spektrum ve frekvenční oblasti. Jednoduše tedy o signálu dokážeme zjistit, jaké harmonické obsahuje . Jeho základem  je symetričnost spektra okolo středu. (Za předpokladu, že se jedná o reálný signál ). Jeho hlavní výhodou je fakt, že k výpočtu nemusíme využívat definiční vztah Diskrétní Furierovy Transformace (Který je příliš náročný pro výpočet). Místo toho si vstupní vzorky rozdělíme na 2 poloviny a každou z nich pak na další 2 poloviny … atd. (Je tedy nutný počet vstupních vzorků    2 na x :  2, 4, 8, 16, 32 … )Výsledkem je, že dostaneme N/2 Diskrétnéch furierových transformací o délce 2, které postupně opět skládáme do jedné velké.

  • FFT transforms a real sampled signal from the Time domain into the Frequency domain. The core of FFT lies in the fact,that a spectrum of a real signal is symmetric around the zero frequency. So we dont need to calculate DFT with its definition,instead we create 2 half-length sequences which are divided into another 2 half-lenght sequences and so on till we have lots of DFT with a length of 2. They are calculated with a so called Butterfly operation and are in further steps joined together to create the whole spectrum of the default signal.

Vzhledem k tomu, že mi přijde zbytečné sem kopírovat obsah celé práce, tak sem pro zájemce dám odkaz na stažení ve formátu .docx a .pdf .

I dont think its necessary to copy the whole document here onto the web, so feel free to download the .docx or .pdf format.

 

Na ukázku je v materiálech spočítaný jeden příklad na metodu FFT o N = 8. Sám totiž mám nejraději, když si mohu ověřit, že daný algoritmus skutečně funguje a když vidím, že výsledky na papíře jsou téměř stejné jako v příkazovém řádku. K slovu ,,Téměř” : GPU má obrovskou výpočetní sílu a pokud používáme datový typ float (Single Precision), jednotlivé matematické operace zanášejí do výpočtu chybu, která by se sice dala odstranit použitím double, ale vesměs se to nedělá a mě to přijde zbytečné. Člověk se akorát nesmí nechat rozhodit tím, že double ukazuje 4,6 a float 4.5 … S tím prostě nic neuděláme ;) Z hlediska paralelní platformy je pak možná dobré se zaměřit na přemapování dat a jednotlivé ukazatele. Já sice vím, že používání funkce ,,mod” je časově náročné, ale na druhou stranu je to jedna z nejlepších metod, jak ,,algoritmu říct, co kdy kde a jak má počítat”.

  • As an exemple,there is a step-by step solution of calculation a simple 8-samples long FFT. As i like to compare the results with a paper version. Please note that when using a GPU for calculations,there might be some error regarding the fact,that we are using single-precision arithmetics. Furthermore i have been using the ,,mod” operator inside the code,which is a slow operation,however really helps to remap running threads onto their correct position.

Dále bych doporučil přečíst si něco o architektuře GPU. Jaké používá paměti a jaké jsou možnosti synchronizace mezi vlákny. To byl totiž pravděpodobně jeden z nějvětších problémů při implementaci algoritmu. Měl jsem na výběr: Búď implementovat rychlejší verzi s omezeným počtem vstupních prvků na 512 a nebo pomalejší verzi se vstupním počtem prvků 2 na 18 = 262 144. U obou algoritmů přitom nedoporučuji tato čísla zvedat ač oficiálně by měly oba fungovat až do (1024 a 1024 x 1024 ). Nakonec jsem se rozhodl implementovat obě dvě a ještě navrch CPU Verzi :) Samotný přepis algoritmu z CUDA do OpenCl je opět v materiálech a bude vysvětlen na obrázku.

  • If you are familiar with a GPU architecture, you already know that there are some blocks running independently and that the only way to sync your algorithm is to launch a new kernel in case your running multiple blocks or to use __syncthreads(). This however works only for sequences of up to 2048 samples (The original version works only for 1024 or 512 samples depending on the compute capability of the hardware). The multiple kernel launching may in fact be faster than the sync inside blocks version. This is caused by the simple fact,that GPU computations became usefull when they have enough work to do. You can see the result of comparing a parallel and serial implementation  of the FFT.

Porovnání výkonu

PCI-EPure

 

Jak je vidět z obrázku, tak přibližně do vstupní délky 16384 vzorků je rychlejší CPU, které nemusí kopírovat data po sběrnici, ani čekat na spuštění Kernelu. Nicméně výpočet řeší sekvenčně a s narustajícím počtem prvků roste náročnost kvadraticky, zatímco v případě OpenCl a CUDA roste Lineárně. Pro délku 2 na 18 jsem naměřil přibližně 9x násobné zrychlení oproti CPU. (Nutno podotknout, že kódy NEJSOU optimalizované). Na prvním z obrázků je rychlost výpočtu se vším všudy, tedy v případě CUDA a OpenCL i s časovou prodlevou kopírování a spuštění kernelů. V případě OpenCL je pak pravděpodobně vidět i čtení kernelu ze souboru, což je odstraněno v následujím obrázku, kde je porovnáván jen čistý výpočetní výkon.

  • As you can see,the CPU is faster then the GPU till we reach about 16384 samples. Since that amount of work, the GPU can compute much faster. The initial handicap is caused by data copying  into the GPU memory over the PCI-Express bus and back to the CPU when the work is done. I have measured the execution time of the GPU to be around 10 times faster than the CPU version when we use 262144 samples. This is a good performance improvement, but in fact I know that it can be much higher. The Code i wrote is simply not optimized. The first graph shows the complete computation time while the second shows a pure performance of the GPU and CPU without data copying,initialization and so on. 

Veškeré Zdrojové Kódy jsou pak k dispozici zde: