Încercarea numerelor cu CUDA (Partea 2: La luptă!)

La numere mă mai gândeam să fac în așa fel încât, pe lângă reducerea câmpului de operare la numerele strict necesare, să optimizez mai mult lucrul cu ele, să mai scadă numărul de operații de înmulțire. Formula deșteaptă de căutare a numerelor prime din cudaSieve, alături de o marcare specială a pattern-ului de numere prime de la countPrimes, îmi dăduse idei să mă mai gândesc cum să lucrez mai direct cu puterile de factori primi, anume să stochez direct aceste puteri de factori primi în fișiere dedicate, fiecare grup de fișiere fiind alocat câte unei plaje de numere prime (pentru numerele de maxim 52, de cifre, pentru cele de maxim 60, de maxim 70, 80, 90... până la 1910, dar fără a lua chiar totul din 10 în 10), iar la construirea sumei de divizori a numărului doar să verific direct divizibilitatea numărului cu acele puteri din fișiere, și cu salt din 2 în 2 indecși de putere pentru o mai rapidă parcurgere, dar fără înmulțiri repetate pentru construcția unei puteri temporare de factor prim, cu care să verific divizibilitatea numărului. Așa că mai scădea numărul de operații aritmetice (o procedură importantă în analiza și optimizarea algoritmilor, și studenții IT-iști învață ceva despre aceasta). O dată scăzut numărul de numere de pornire la cel strict necesar, puteam în schimb mări cantitatea de coeficienți de legătură dintre numere, despre care mai vorbeam, adică alte numere bazate pe produse de puteri de factori primi componenți.

Dar până să ajung acolo cu reorganizarea sistemului de fișiere numerice, au fost destule zile de lucru și m-am mai ocupat și cu alte lucruri legate de hardware și de software la numere. Inițial știam că tehnologia CUDA se folosește eminamente cu limbajul C (și cu forma sa de OOP, C++, pe c are însă nu o voiam ca să nu complic treaba cu overhead-ul generat de clase). În primele zile din noiembrie 2022, am aflat că și limbajul Julia (unul mai tânăr) folosește CUDA, iar în primul sfârșit de săptămână din luna noiembrie am făcut experimente pe el, instalându-l pe calculator (încă stăteam pe SLI PLUS) și testând funcționalitatea tipurilor întregi pe el. Văzusem că are și două tipuri întregi pe 128 de biți, unul cu semn și altul fără semn (UInt128), însă cel fără semn nu permitea decât reprezentare hexa, cu 0x și multe cifre (maxim 16, adică 16 octeți ca să ajungă la 2 la puterea 16*8 = 128), așa că pentru reprezentarea cifristă îmi rămânea tipul Int128, care chiar și așa tot ajungea pe la 39 de cifre zecimale, trecând de 10 la puterea 38. Pe Julia nu prea mi-a plăcut faptul că, dincolo de acomodarea inițială cu instalarea a tot felul de pachete (inclusiv unul cu nume ciudat, Cthulhu), dacă ai o eroare îți pune o veritabilă diaree programatoristică pe ecran, cu un mesaj luuung cu eroarea. Iar când am încercat să folosesc pe CUDA acel tip de date cu Int128 și să fac împărțire, am remarcat că... nu este implementată divizibilitatea pentru CUDA cu acel tip întreg de date! Dădea eroare la compilare. Pentru CPU-ul normal se puteau face operații cu întregi de 128 de biți și la împărțire, dar tot fără a putea reprezenta cu cifre normale numerele fără semn. Nu puteam renunța complet la împărțire pentru lucrul cu numerele, cu sau fără CUDA. Și încă un lucru - deși citisem cum că Julia poate fi mai rapidă decât C, la testarea unor operații cu întregi elementari pe CUDA am văzut că Julia se mișcă mai încet decât C, deci după 6 noiembrie 2022 am lăsat-o deoparte (măcar am încercat și-am văzut ce înseamnă Julia!) și m-am orientat înapoi spre C.

Deși primele impresii cu CGBN-ul nu fuseseră foarte convingătoare, am remarcat că, până la data aceea, biblioteca aceasta de numere naturale mari pentru CUDA este singura care se apropie de ceea ce aveam eu nevoie să fac pe video, având operații aritmetice care să abordeze nevoile mele numerice - cu acomodarea inițială la noua sintaxă, CGBN având și comunicare cu GMP pentru convertirea numerelor obținute pe video la un format reprezentabil în scris cu ajutorul GMP-ului. Căci în funcțiile de video puteam face printf pe CUDA numai cu numere întregi clasice, suportate direct în CPU, dacă încercam ceva cu formatter de GMP nu mă lăsa, iar pentru tipurile de CGBN nici măcar  nu exista un formatter specific. Am văzut că pe CGBN numerele întregi sunt reprezentate prin tipul cgbn_t, am învățat cum să le instanțiez (nu se putea decât pe video, cu declararea typedef-urilor pentru mediul de cgbn_t la începutul fișierului cu codul-sursă, o vreme am făcut typedef-urile într-un header separat, dar niște erori m-au determinat să le pun în fișierul .cu direct, și au fost zeci de astfel de fișiere făcute de-a lungul lunilor, până în februarioe 2023) și am remarcat că și numerele mari cgbn_t se pot combina aritmetic cu întregi clasici, în operații de adunare, scădere, înmulțire, împărțire, divizibilitate, shiftare binară, rest..., dar că spre deosebire de situația de pe CPU cu GMP-ul, unde era voie cu întregi fără semn de 64 de biți, pe video la CUDA cu CGBN se acceptă numai numerele pe 32 de biți, adică până la 4294967295.

Și atunci, în prima parte din noiembrie 2022, după dezamăgirea cu Julia, am făcut mai multe încercări de a convinge codul CGBN-ului, descărcat la mine pe calculator, să suporte și aritmetica cu numere native de 64 de biți. Am reușit parțial pentru adunare și înmulțire (doar pe bitul inferior, la o primă vedere operațiile mergeau bine dacă foloseam întregi nativi peste 4294267295 - în caz contrar 4294967296 devenind 0 și făcându-se astfel operații modulo 4294967296 cu numerele native mai mari de pragul de mai sus, ajungându-se și la împărțirea cu 0), iar la scădere și mai ales la împărțire, la care nu puteam renunța de tot, rezultatele erau razna. Și încă nu ajunsesem la shiftarea cu peste 32 de biți.

Am umblat prin măruntaiele codului-sursă al CGBN-ului, inclusiv la instrucțiunile de assembler pe CUDA (nou pentru mine, diferit de assemnbler-ul clasic de C pe care îl cunoscusem în vara lui 2010, și numit PTX Assembler), am văzut acolo cum sunt definite operațiile pe regiștri de 16, 32 și 64 de biți, la operațiile adăugătoare precum adunarea și înmulțirea puteam forța nota cu 64 de biți într-adevăr, dar la împărțire erau niște instrucțiuni speciale cu clamp și cwrap care nu erau recunoscute pentru regiștrii pe 64 de biți, ci pentru 32, și acolo mi-a fost clar că este o punte de netrecut. Mai explorând codul am văzut și că reușita mea cu 64 de biți, acolo pe unde era, se aplica doar la low bit și nu și la high bit, ca să fie treaba completă. Mai era oricum de verificat și la shiftări; am intrat și prin logica de multi-threading, unde se făcea shuffling și sincronizare între threaduri pentru lucrul cu instanțele sau iterațiile de lucru pe video (numite instances), iar CGBN-ul are un număr reglabil de thread-uri per instanță, numit TPI (Threads per instance), care este o putere de 2, de obicei 4, 8, 16 sau 32.

Și pentru că se lucrează cu calupuri de câte 32 de biți pe video, fiecare număr mare din cgbn_t este împărțit pe limb-uri de câte 32 de biți (pe care eu încercam să le schimb în 64, cu efecte asupra numărului și structurării lor în codul CGBN-ului), iar în funcție de mărimea în biți setată ca precizie, variază și numărul de limb-uri - de pildă dacă alegem constanta BITS să fie 256, înseamnă că lucrăm cu numere cgbn_t de până la 2 la puterea 256, fiecare număr are 8 limburi de câte 32 de biți, poate să le umple pe toate sau  nu, și TPI-ul poate să fie de la 4 la 32 (eu am umblat și pe-acolo, încercând să fac TPI de 64 de exemplu). Aveam să remarc pe parcurs că CGBN-ul este incomplet, în sensul că nu are configurarea pusă la punct pentru când numărul de limburi este mai mare decât TPI-ul, de pildă dacă TPI este 32 și numărul de limburi este 64, iar BITS este 64*32 = 2048 - sau când BITS este și mai mare. Prin ianuarie 2023 m-am izbit de eroarea cu Incomplete type not allowed, deoarece lipsea configurașția pentru is_multi, atunci când LIMBS este multiplu de TPI. Pot să încerc să rezolv local problema aceasta, dar îmi trebuie mo bună înțelegere a definițiilor de acolo ca să reușesc. În specificațiile CGBN-ului, așa puține câte sunt ele, că nici nu este o bibliotecă mare, scrie că se suportă numere de până la 32K biți (32768 cred că vrea să spună), care pot fi reglate din 32 în 32 de biți, deși mai simplu este ca precizia dată de constanta BITS să fie o putere de 2, între 32 și 32768.

Iar TPI-ul este aranjat să suporte valorile 4, 8, 16 și 32. Pentru situația când BITS ar fi 32768, LIMBS ar fi 1024, adică de 32 de ori mai mare decât maximum de TPI, și atunci configurația lipsă cu is_multi ar trebui să acopere situațiile de până la LIMBS = 32*TPI, sau chiar 256*TPI când TPI este mai mic - deși cu cât numerele sunt mai mari, un TPI mai mare ar ajuta mai repede la facerea numerelor, adică mai multe threaduri deodată la luptă pentru reprezentarea unui ditamai număroiul, deși atunci nici viteza de lucru nu ar fi cine știe ce. Să văd cu configurația de dlimbs_algs_multi pe care cei de la NVIDIA Labs nu au făcut-o încă!

Anul trecut când vedeam că tot nu merge să conving CGBN-ul să suporte aritmetica nativă pe 64 de biți pentru CUDA, la fel ca la CPU-urile obișnuite ale vremii, am citit că pe plăcile video lucrul pe 32 de biți cu numerele este normativ, spre deosebire de CPU-uri, că este într-adevăr o altă lume. Atunci mi-am zis că oamenii de la NVIDIA Labs care au făcut CGBN-ul știau ei ce fac atunci când au pus la punct aritmetica pe 32 de biți, și că nu are rost să încerc să mă dau eu mai deștept decât ei cu cei 64 de biți, așa că am luat iarăși codul-sursă original al CGBN-ului și nu am mai lucrat cu suprascrieri forțate de logică pentru 64 de biți. Pe-acolo prin cod sunt tot felul de directive complicate pentru înpărțirea muncii între threaduri (câteodată eu schimbam măștile de sincronizare pentru threaduri, care trebuiau într-adevăr să fie pe 32 de biți, cu valori pe 64 de biți, adăugând opt litere de f la cele opt litere f deja existente în mască. Chiar că nu avea sens). Și mai sunt shiftări printre biții mai semnificativi, mai puțin semnificativi... nici nu mai țin minte bine. Poți să adaugi zerouri în stânga (leading zeros) sau în dreapta (trailing zeros).

Și am început să pun treptat la punct un sistem de algoritmi optimizați și fișiere de lucru cu numerele, cu headere noi și file cu extensia .cu, cu CGBN-ul pe 32 de biți, pentru numere. A fost o muncă treptată, pe măsură ce înțelegeam cum să convertesc codul de la GMP la CUDA și cum să fac alocările specifice și eliberările de memorie pentru video, preluarea de date numerice către și dinspre placa video, convertirea rezultatelor într-un format acceptabil GMP, folosirea adecvată a funcțiilor aritmetice cu cele 4 operații și împărțirea cu rest, parametrizarea corectă a funcțiilor în funcție de ceea ce trebuia să facă algoritmul, curățarea parametrilor inutili, modificarea în timp scurt a tuturor fișierelor, prin folosirea deșteaptă a directivei Replace prin mai multe fișiere deschise și grija de a impacta numai ceea ce trebuie; am observat comenzile care erau incompatibile cu placa video, și care în device mode dădeau eroare că nu sunt definite în device mode; și dacă puneam __host__ și __device__ în fața aceleiași funcții de video, eroarea la directivele care nu erau pentru video (de ex. operații cu numere de GMP) nu mai apărea, dar ele erau pur și simplu ignorate. Așadar, când suntem pe CUDA cu CGBN, nu are rost să punem mepezetele prin funcțiile alea. Există în codul pentru CPU (host) instrucțiuni speciale pentru ca mepezetelele să fie preluate de video la CGBN, iar apoi rezultatele venite de pe video să fie convertite înapoi la GMP ca să fie citite. Nici măcar operațiile cu stringuri la care să fie convertite numerele nu se acceptă în device mode. Am dus destulă muncă de lămurire pentru eliminarea de variabile și parametri inutili și redenumirea unor tipuri de date cu aliasuri, pentru economie de cod. Dar cu aceste lucruri depășim puternic cadrul calendaristic al lunii noiembrie 2022!

Iar în primă fază, am încercat să generez numere prin randomizare, experimentând numerele aleatorii pe CUDA fie prin definirea lor în același nucleu (kernel function, cu __global__) cu cel în care voiam să generez numerele (mi-am stricat rezultatele făcând astfel), fie să generez semințele aleatorii în  mod separat și să lucrez cu rezultatele aleatorizate pe kernel-ul propriu-zis pentru numere. Și am  mai căutat o soluție alambicată în care să stabilesc dacă un număr respectă definiția matematică fără a mai fi nevoie de operația de împărțire, dar se genera overhead. Împărțirea făcea fițe pe CGBN, uneori nu genera rezultate, nu se făcea nimic cu numerele. Mai târziu aveam să remarc că primeam, printre rezultatele numerice corecte de la CUDA, și niște numere parazite, care nu respectau abundențele cerute. CGBN avea să mă dezamăgească cu imperfecțiunile sale în iarnă, și nici câștigul de viteză nu era unul exponențial - față de procesorul 7950X, la numerele de 256 de biți, poate că era o mărire de viteză de 1.5-2x (plus posibilitatea de a rula 3-4 programe în paralel), dar la numerele mai mari avantajul scădea, iar pentru buba de la LIMBS > TPI, nu mai funcționa deloc CGBN-ul, trebuia să mă întorc la GMP. Deci a fost o idee foarte bună să fac și pentru GMP algoritmi noi, cu ei aveam să-mi aduc majoritatea rezultatelor numerice noi, în a doua parte a iernii. AJungem și acolo!

Începând să mă ocup de migrarea numerelor pe CUDA în limitele și sub auspiciile CGBN-ului în noiembrie 2022, începusem să mă preocup mai serios de o înnoire a puterilor video din casă. Spre deosebire de luna aprilie 2019, aveam acum și niște bani ca să-mi împuternicesc dotarea video, ce nu făcusem până în 2019. Aveam și CGBN-ul la îndemână, spre deosebire de bibliotecile nemulțumitoare găbjite atunci pe internet - precum CUMP, CAMPARY și ce altceva o mai fi fost. Știam despre noile procesoare AMD și Intel, cu suport DDR5 și PCI-E 5, plus noile plăci video bestiale marca 4090, apărute primele, urmate de 4080 și-apoi și de restul suitei 4000, cu 4070, 4060 și TI-urile. Ar mai fi fost și 4090 Ti, dar aceasta ardea pur și simplu, și au trebuit să îi dea discard cei de la NVIDIA. La 4090-uri am observat că au 16384 de nuclee CUDA, care pot depăși 2500-2600 de MHz în frecvență la plăcile cele mai bune (și mai încolo aveam să văd că prin overclocking poți să bați chiar și cota 3000, apropiindu-te mai mult de frecvențele obișnuite ale threadurilor de CPU). Neștiind încă cum avea să se miște CGBN la mine la numere, credeam că 16384 de threaduri de vreo 2.6 GHz fiecare ar trebui să dea rezultate mult mai rapide decât 128 de threaduri de 4.5 GHz fiecare (4090 top versus Threadripper top), pe lângă faptul că plăcile video, spre deosebire de CPU-uri, sunt plămădite înadins pentru operații cu numerele! Cum ar fi să compari o forță totală de vreo 42600 de GHz cu una de 576 de GHz, și la un preț de vreo 3 ori  mai mic pentru cea cu teraherții (4090 versus 5995WX Threadripper), plus factorul de „înadins pentru numere” versus „nu neapărat pentru numere” din confruntarea GPU cu CPU! Dacă lucrurile ar fi stat așa de strict aritmetic, ar fi trebuit atunci ca o placă video 4090 să-mi dea o forță cât vreo 75 de Threadrippere din alea pentru vreo 30% din preț. Sau cât vreo 240 de procesoare AMD 7950X, nou lansate atunci, toamna! Și fără să puem în calcul diferența generațională și de litografie, că nu contează numai numărul de threaduri și frecvența lor.

Uitându-mă prin noiembrie 2022 la posibilele plăci video noi, și știind că vine și Black Friday, îmi măream treptat pretențiile pentru numărul de nuclee CUDA pe care să-l aibă placa. Dar noile plăci video suportau PCI-E 4 (nu 5, ce-i drept), și eu în casă aveam PCI-E 3, chiar și pe vechiul Threadripper 2990WX, și plăcile video cu generația asta nu ajungeau prea sus cu CUDA Cores. Așa că după povestea cu visul din octombrie, am început să pregătesc mai serios un upgrade la DDR5, PCI-E 5 (compatibil cu 4, firește) și un CPU de maxim 32 de threaduri, la preț mult mai mic decât un Threadripper de top, dând în schimb loc unei plăci video 4090, ca să fie progresul pe cinste și să nu mă mai complic cu plăci intermediare. Chiar și CPU-ul avea să fie de top pe segmentul lui!

Pe parcursul lunii, încă mă mai gândeam că poate prin 2023 o să-mi fac calculator nou, nu chiar atunci, în 2022; și că noul calculator va putea fi cu Threadripper și DDR5, dar cu un CPU mai mic, la 24 de nuclee și 48 de threaduri, nu un flagship. Și că în 2022 mai pot să-mi iau o placă video de PCI-E 3.0. Dar treptat am aspirat mai sus! După mijlocul lui noiembrie 2022 eram mai hotărât că s-a apropiat timpul pentru DDR5 și placă video puternică, și fără Threadripper. Nu mai era nevoie de investit într-un mare CPU.

Și am pus ochii pe o placă video 4090 de la MSI, cu radiator de răcire lichidă, MSI Suprim Liquid 4090, care atingea 2640 MHz în boost. Am ales AMD să-mi fie configurație, cu suport strict pentru DDR5, cu un număr mai bun de magistrale PCI-E 5 (24, față de 20 la Inteluri), procesorul 7950X cu 32 de threaduri și frecvență de 4.5 GHz (5.7 la boost), și o placă de bază beton, MSI MEG X670E ACE, să nu fie chiar GODLIKE, al cărei nume mi s-a părut impios, și ceva mai ieftină să fie placa, dar tot cu calitate: suport pentru SSD PCI-E 5 NVME, 10 Gbps la internet, mai multe sloturi de PCI-E 4 și loc de două plăci video PCI-E 4 cu regim full (x8 la PCI-E 5 este ca x16 la PCI-E 4).

Dar nu mă hotăram odată să dau drumul la comenzi. Așteptam o scânteie motivațională ca să dau drumul la comenzile de componente pentru noul calculator. Mă gândeam că aș putea refolosi o carcasă din casă. Eram în concediu în a doua jumătate a lunii noiembrie 2022 și mai meșteream pe la noile programe cu CUDA, îmi alesesem deja lista precisă de componente dacă ar fi să fac un nou calculator, inclusiv pentru o nouă carcasă mă uitasem, dar până atunci am remarcat că în casă mai puteam să-mi îmbunătățesc configurația funcțională - pe SLI PLUS placa 1050 Ti era forțată să stea în regim x8, la jumătate din potențial, dar pe ASUSPRIME sau Threadripper (nu voiam să mai implic XPOWER-ul și Aorusul) putea sta cu x16, la maximum de potențial.

Miercuri, 23 noiembrie 2022, am vrut să mut la ASUSPRIME placa 1050 Ti, să mă mut de pe SLI PLUS pe ASUSPRIME ca să nu angrenez ditamai Threadripperul la consum în poveste. Dar când am vrut să dau drumul la sistem pe ASUSPRIME, tot aveam o eroare la POST cu un device overpowered sau ceva de genul - ca și când un component atașat la calculator ar fi avut mai mult curent decât trebuie. Ceva din familia „Power surge on the USB port”. Nu reușeam să-i dau de cap; de-a lungul zilei am încercat în fel și chip, aveam dificultăți la decuplarea plăcii video 1050 Ti, care era dual-slot, atunci când trebuia să o scot de pe o placă de bază; văzusem între timp, în noiembrie, pe când îmi alegeam noile componente, că există și Riser Cables care te ajută să pui plăcile video mai departe de placa de bază, lăsând astfel mai mult spațiu pe placa de bază - mai ales la cablurile mai lungi - indiferent de cât de multi-slot este placa vide, și făcând mai ușoară și cuplarea/decuplarea. Cu placa mea video m-am plimbat între ASUSPRIME, SLI PLUS și Threadripper, în timp ce Asuspraimul tot nu-și revenea din eroarea cu device overpowered; la Threadripper placa video refuza să pornească bine (altă problemă, faptul că avea alimentare separată PCI-E nu era tolerat de sursa de 1600 W de acolo, cu sau fără Eco Mode, ciudat lucru), indiferent de cum configuram eu sursa, să fie pe economic mode sau nu - nu voia să o pornească, mergea doar placa video de mai mică putere, de 2 GB, 1050 simplu. Și fără alimentare separată.

Și seara, când eram la un moment dat pe SLI PLUS cu placa video de 4 GB de la GIGABYTE, voind s-o scot și neavând răbdare să dibuiesc cu unghia clama de desprindere pe sub placă, din slotul PCI-E 3, și deci tot trăgând de ea, am smuls-o cu tot cu suportul de plastic pentru acel slot de PCI-EXPRESS al plăcii SLI PLUS, mutilând astfel placa de bază (slotul principal de video/PCI-E devenit nefuncțional). Iar pe urmă, la THREADRIPPER și Veriton, am văzut că placa video GIGABYTE nici măcar nu mai pornea. Nu i se mai învârtea ventilatorul, nu se mai în călzea, necum imagine pe ecran, nimic.

Mutilasem o placă de bază și stricasem placa video cea mai bună din casă, după o eroare stupidă de power surge la ASUSPRIME, când voisem să îmbunătățesc funcționarea configurației video din casă. Rămăsesem și cu o placă de bază demantelată (cea mai mică dintre cele mai bune), pe care și așa o mai înlocuisem o dată în 2018, când începusem creditarea la bancă, și mai aveam și o placă video în minus - tocmai pe cea mai bună dintre cele vechi! Era un mic dezastru în penultima miercure din noiembrie.

Aceasta a fost scânteia de motivație pe care o așteptam ca să mă hotărăsc odată să iau componente noi pentru calculator!