Încercarea numerelor cu CUDA (Partea 1: Prevestirea)

Joi, 27 octombrie în 2022, de dimineață, am avut un vis. Se făcea că era dimineața de 5 ianuarie 2023, tot o zi de joi, cu cer senin așa de dimineață, și eu tocmai primisem niște componente noi de calculator, în special un procesor nou, care era într-o cutie de culoare albastră. Știam pe atunci că se lansează niște procesoare noi de la AMD și Intel pentru consumatori, cu până la 32 de threaduri (procesoare logice), iar în vis mi se încâlcea ideea că acela este fie un CPU din acesta, fie un Threadripper cu 128 de threaduri care însă nu putea să facă parte decât dintr-o generație m ai veche, cu DDR4, noile CPU-uri de mai sus fiind deja și cu suport DDR5.

Și în plus, tindeam în vis să mă gândesc că mi-am luat un CPU din acela cu mai puține threaduri, nu doar pentru progresul memoriei, ci și pentru că, în perspectivă, mă gândeam ca la o viitoare încercare cu numerele să nu mă mai concentrez pe CPU-uri foarte puternice precum Threadripper, ci pe plăci video foarte puternice, cu mii de procesoare logice, și cu CUDA, crezând că forța lor de procesare mă va ajuta să obțin rezultate numerice cu o viteză exponențial mai mare decât pe un CPU obișnuit. Și în orice caz, mă cam liniștisem de ceva timp cu marile cheltuieli de hardware - abia ce mai luasem câte ceva prin toamna lui 2022, întâi pentru calculatorul cel mai vechi din casă, în septembrie, apoi mă mai dotasem și cu o cameră web nouă pe 4K și echipamente router+cablu de 1/10 Gbps, cu ocazia upgrade-ului Internetului din casă de la 100 Mbps la 1 Gbps, în octombrie: luni pe 3 fusese instalarea, marți pe 4 cumpărasem o cheie Retail pentru Windows 11, de pe site-ul eMAG, ca să pot clona și muta sistemul Windows 11, cu DiskGenius, pe alt calculator mai puternic decât Veritonul și pe un mijloc de stocare mai rapid, ca NVME-ul. Că îl aveam până atunci pe un SSD SATA din 2015. În primă fază l-am mutat pe calculatorul cu SLI PLUS, dar eroarea cu sistemul care nu întrunește condițiile necesare pentru Windows 11 m-a urmărit și acolo, trebuia să schimb niște parametri într-un fișier DLL și în configurația de regiștri a sistemului ca să mă lase, la fel ca pe Veriton.

Abia pe ASUSPRIME puteam scăpa de problemă, căci procesorul i9-7900X a fost la urmă acceptat (inițial nici el, în 2021), alături de Threadripper 2990WX. Pe SLI PLUS aveam i7-7740X. Iar pe Veriton, i5-650X de generație 2010. Mai luasem și un NVME de jumătate de terabait în toamna lui 2022, cu generația PCI-E 3.0, tot în octombrie, ca să mut pe el Windows-ul clonat cu DiskGenius, și în primă fază rămăsesem pe SLI PLUS cu el, dând astfel pentru o vreme ocazia calculatorului respectiv să fie refolosit, căci din 2019 de când mă potolisem cu numerele, în afara perioadei mai-iulie 2021, sătăusem în principal numai pe Veriton, cel mai mic și mai ieftin calculator din casă, mai deschisesem Threadripperul, dar nu mult, și pe celelalte ocazional, pentru câte un update de BIOS, în măsura în care au mai apărut. Consumul de kilowați era și el potolit de mult la două cifre lunar, uneori chiar pe la 30 de kilowați sau chiar mai jos, mai crescuse în cele câteva luni de coding numeric din 2021, dar nu mai sus de 91 pe lună, cu ajutorul lui Threadripper - și acum, în toamna lui 2022, SLI PLUS-ul mai împingea și el ceva kilowați în plus la factura lunară, că nu consuma la fel de puțin ca Veritonul, dar tot fără să mă reped tare în sus cu consumul. Nu aveam în plan un upgrade hardware major pentru 2022, dar în perspectivă mă gândeam că prin 2023 s-ar putea, eventual cu o generație Threadripper cu DDR5, dar fără a ține neapărat să iau cel mai mare CPU posibil din ea, ci cu accent pe partea de CUDA, asta dacă aveam să găsesc o bibliotecă software pentru numerele întregi mari pe CUDA, fapt nerezolvat în 2019, când aflasem despre puterile CUDA și nu găsisem ce căutam (și oricum, nici nu puteam să-mi mai iau atunci ceva bun pe video, cu cheltuielile care fuseseră).

Dacă nu s-ar fi rezolvat problema cu biblioteca de CUDA, știam că m-ar costa mai mult investiția într-un Threadripper de top, decât într-o placă video cu maximum de nuclee CUDA, dar nu voiam să mai iau ceva pe DDR4. Iar Threadripperele sigur mai aveau de așteptat până acolo, sigur nu pentru 2022. Voiam ca următorul calculator din casă, dacă mai avea să fie, să aibă DDR5 și suport PCI-E 5, lucru care nu era deocamdată posibil decât pe noile configurații cu CPU-uri nu prea bogate. Și atunci, ca să și merite așa ceva, musai să fie și CUDA cu video, ca să am ceea ce credeam că va fi un câștig major de viteză la lucrul cu numerele. Visul de mai sus m-a pus puțin pe gânduri, că nu voiam să-mi mai iau ceva nou chiar în 2022. Iar dimineața de 5 ianuarie era, ce-i drept, după Anul Nou în 2023, suna interesant, în vis singurul lucru cât de cât clar era procesorul cu cutie albastră, proaspăt livrat, mai era ceva cale de parcurs până la asamblare, ideea era că am niște componente noi.

Am zis pe moment că este un vis, că tot nu plănuiesc să fac ceva de genul acesta, dar mi-a pătruns treptat la suflet ceea ce am visat. În a doua parte a anului 2022, nu țin minte data exactă, găsisem pe Internet o bibliotecă de numere mari pentru CUDA pe care nu o văzusem în 2019, unde erau descrise câștigurile de viteză în funcție de mărimea operanzilor numerici (mărimi exprimate în biți, pe puteri de 2, ca 2048 bits, 4096 bits...) și categorisite pe diferite operații aritmetice, ca adunarea, scăderea, înmulțirea, împărțirea, cel mai mare divizor comun și nu numai, să se vadă de câte ori merge mai repede pe o placă video Testa V100 (Volta) decât pe un CPU Intel Xeon cu 20 de nuclee, numit E5-2698v4. Biblioteca se cheamă CGBN,  sau Cooperative Groups Big Numbers, unde grupuri de threaduri video lucrează împreună ca să reprezinte numere întregi cu multe cifre și să proceseze operații aritmetice cu ele. Motivația numărului de thread-uri puse împreună este faptul că acele numere au nevoie de un efort mare de procesare, deși pe GMP la CPU-uri nu am remarcat o astfel de împărțire a trebii pe threading. Și deci o placă video Tesla versus un Xeon!

Placa asta video are 640 de Tensor Cores, 16 sau 32 de GB de memorie de tip HBM2 (dar pe repository-ul public de Github al bibliotecii CUDA nu scrie de care V100 din asta a fost), în funcție de dacă este pentru PCI-Express sau NVLink, și pare a avea 5120 de nuclee CUDA, deși pe site-ul oficial al NVidia nu scrie câte are. Și frecvență de la 1245 la 1380 de MHz. Iar procesorul acela de Xeon are patruzeci de procesoare logice (Thread-uri), poate să ajungă la 3.6 GHz (neoverclockabil). Sporurile de viteză pentru placa video scădeau cu creșterea numărului de biți, de la 128 de biți la 8192, și se câștiga cel mai mult la adunare și scădere, înmulțirea era undeva la mijloc, iar împrățirea mai jos, oricum este o operație mai grea. Pentru 128 de biți lucrurile se mișcau de peste 174 de ori mai repede la adunare și de 30 de ori mai bine la împărțire, iar la 8192, de 28.6 ori la adunare și 7.8 la împărțire. Și se făcea și o medie!

Putem să vedem ce au făcut oamenii: https://github.com/NVlabs/CGBN

Nu am explorat atunci în amănunt detaliile acestea ale specificațiilor comparative pentru procesor și placa video, și ce-i drept nici nu mi s-a părut prea convingător CGBN-ul, cu sporurile variate de viteză, precizia limitată în exprimare (cu salturi exponențiale de puteri de 2 la numărul de biți, pe câtă vreme la GMP/MPIR și MIRACL precizia era automată sau reglabilă după placul utilizatorului) și tendința de a scădea mult la numerele mai mari. De fapt, preciziile lor se pot da din 32 în 32 de biți, însă sunt și lucruri incomplete și imperfecte prin bibliotecă, mai trebuie lansată o versiune a ei îmbunătățită, o să ajung și acolo. Atunci în 2022, nu mi s-a părut cine știe ce de la început CGBN-ul, dar măcar aducea un suflu nou față de neîmplinirea de la cercetările din 2019. Că și în 2019 sigur aș fi băgat CGBN-ul de seamă dacă l-aș fi văzut. Pare a fi fost lansat chiar în acel an, dar în primăvară nu l-am văzut, poate că l-am pierdut la o lună-două distanță.

Începând din 31 octombrie 2022, am pornit preocuparea mai îndeaproape pentru algoritmii de CUDA. Mai rulasem ceva program cu numere zecimale mici pe CUDA, pe una dintre plăcile video din casă, și în toamna lui 2022 mi-am recapitulat iar dotarea video de acasă. Îmi mai luasem în septembrie 2022 o placă video NVIDIA GT 730 pentru Veriton, ca să scap de vechitura cu 9500 GT, dar tot slabă era (GIGABYTE GeForce GT 730 2GB GDDR5 64-bit), la 902 MHz, 1253 MHz clock de memorie și 384 de nuclee CUDA. Aveam în total cel puțin 2 astfel de plăci și 2-3 GT 1030-uri, tot cu GDDR5 de 2 GB, și tot cu câte 384 de Cude pe ele și frecvențe de la 1128 la 1468 MHz, la Threadripper aveam o GT1050 de 2 GB cu 640 de Cuda Cores și frecvențe de la 1392 la 1531 MHz, cu 7 Gbps memorie efectivă la clock, iar placa video cea mai bună din casă era NVIDIA GeForce 1050 Ti WindForce (OC), cu 4 GB, 768 de Cuda Cores și ea, 128-bit GDDR5 și 1752 MHz memory clock (înmulțit cu 4, dă 7.08 Gbps), cu frecvențe la nuclee de la 1291 la 1392 MHz (base versus boost). Aceste plăci cu 1050 respectau arhitectura video Pascal.

Cu placa video 1050 TI puteam lucra în regim de PCI-E 3.0 x16, dacă o instalam cu bine pe o placă de bază mai dotată. Adică cu un CPU cel puțin ca i9-7900X. O uitasem pe SLI PLUS, unde rula cu x8 cu acel procesor mai slab, i7-7740X.  Pe lângă noile procesoare cu DDR5 și plăcile de bază ce începeau să se lanseze, știam că apare și o generație nouă și puternică de plăci video, 4090, cu muuult mai multă forță procesoristică decât ce aveam eu acasă și cu frecvențe mai mari de 2 GHz, ele putând ajunge chiar și peste 2.5 și 3 GHz, mai ales în boost și overclocking. Aveam să le descopăr treptat. În primă fază, descopeream pas cu pas CUDA, trecând dincolo de testul inițial cu zecimalele. Așa că începând cu lunea de 31 octombrie 2022, am luat la puricat niște exemple de bază cu CUDA în C, începând cu unul numit „Count Primes”, care să numere câte numere prime sunt în intervalul de sub un oarecare număr natural, și cu altul numit Cuda Sieve, care făcea o căutare asemănătoare, aplicând o formulă deșteaptă de rapidizare a căutării numerelor prime, ceva cu 2*[(2*tmp+3)*i +tmp +1] +1 = 4*tmp*i + 6*i + 2*tmp +2+1 = 4xy + 6y + 2x + 3, unde y este nenul și x nul.

Era impresionant cât de repede se puteau mișca aceste programe cu CUDA și numărarea deșteaptă a numerelor prime! Puteam depăși miliardul cu puțină răbdare, semnificativ mai repede decât la alți algoritmi pe CPU cu primele N numere prime. Am făcut cunoștință mai îndeaproape decât în 2019 cu extensia de .cu a programelor cu CUDA, și cu sintaxa de __global__ a funcțiilor care se adresează direct plăcii video în limbajul C. Mai avem __device__ pentru placa video și __host__ pentru comunicarea cu memoria și CPU-ul de pe placa de bază. Și alocările cu copierile de date în memorie sunt diferite de cele pentru RAM-ul legat de CPU, cu cudaMalloc și cudaMemcpy. Iar elib  erarea memoriei pe video, la finalul programului, avea nevoie tot de o sintaxă specifică de cudaFree.

Am văzut atunci că putem avea metode foarte, foarte deștepte ca să obținem rezultate rapide când vrem să aflăm lucruri despre numerele care ne interesează. Inclusiv din inteligența codului-sursă, până la forța hardware. Iar povestea cu numerele, tot până la o nouă emancipare a forței hardware, mai avea și ea oricum nevoie de o nouă deșteptare a algoritmilor, ca să merg mai repede la miezul problemei descoperirii lor, fără risipa de RAM și de terabaiți de stocare din perioada 2017-2021, cu accent pe până în 2019.

Dincolo de exemplele internautice cu numărarea de numere prime, treceam treptat spre programele mele de interes pentru descoperirea numerelor al căror numitor al fracției abundențiale (suma divizorilor supra număr, ca fracție ireductibilă) să fie maxim 10. Am început cu un program numit sumnum.cu, care lua factorii primi ai numerelor de maximum 52 de cifre și parcurgea, la pas precis, numerele dintr-un interval natural (ca de la 6 spre 6000, din 6 în 6), calculându-le suma divizorilor și verificând că respectă condiția de mai sus. Încercam atunci să-mi dau seama și cum se pot folosi eficient Cuda Cores-urile plăcii video 1050 Ti, crezând că ele sunt identice cu thread-urile cu care CUDA cu C lucrează pe plăcile video (THREAD_NUM), adică punând de pildă ceva de genul

#define THREAD_NUM 768

și respectând mai departe conceptul de blocuri de date pe thread (BLOCK_NUM), care trebuia să fie o putere de 2, ca de pildă:

#define BLOCK_NUM 512

La apelul funcției de __global__ care se ocupa pe CUDA de numere, căutam o abordare de tip bidimensional, în care să se țină cont de THREAD_NUM și de BLOCK_NUM, fără a dimensionaliza mai mult situația. Pentru block ID-uri și thread ID-uri, foloseam doar dimensiunea „x” a acestora, ele putând avea și y și z, dar scopul meu era să fac cât mai optim munca cu numerele pe placa video.

Iar legat de consumul de resurse, mă gândeam să nu mai folosesc drept spațiu de căutare decât numerele de fond 1 (cum le spuneam înainte), care erau 31231 la sfârșitul lunii octombrie 2022, deci cele care respectau exact definiția de mai sus, fără alte numere ajutătoare de manevră, care în anii trecuți reclamaseră o groază de TB de spațiu de stocare și peste jumătate de tera de DDR4 RAM. Și mai departe, să respect principiul obișnuit de a aloca o anumită plajă de numere prime componente, în funcție de mărimea numerelor în care făceam căutarea (mai târziu, după optimizări, putând folosi chiar o plajă mai mare de numere prime, poate-poate mai găsesc ceva numere), și fișiere cu coeficienți de legătură posibili între numere, care însă să nu mai ceară nici ele 100-200 de TB, sau chiar mai mult, și-un cârd de RAM. Iar legat de raportul de preț între placa video și CPU, un Threadripper de top poate să fie de 3 ori (sau mai nou chiar de 5 ori) cât o placă video de top, pe când eu credeam că forța de procesare video pe CUDA poate fi de vreo 50-100 de ori mai mare decât la Threadripper pe placa video de top (aveam să mă lămuresc mai încolo cu CGBN-ul ce și cum).