Villámgyors C++-kódkönyvtárat tett közzé a Facebook

Villámgyors C++-kódkönyvtárat tett közzé a Facebook
2012-06-04T07:05:15+02:00
2012-06-10T09:51:31+02:00
2022-10-25T06:20:40+02:00
  • Ha van humorerzeke az amd-nek, akkor csinal egy setcpuid utsitast es problem solved
    Mutasd a teljes hozzászólást!
  • tudod, hova dugjak maguknak ezek fenyeben? :)
    Mutasd a teljes hozzászólást!
  • Mutasd a teljes hozzászólást!
  • Annak az a baja, h nem intel cpu-ra kikapcsol egy csomo optimalizaciot, fuggetlenul attol, h amugy menne-e vagy sem. Amdm van' szal nem.
    Mutasd a teljes hozzászólást!
  • Kicsit off, de ti használjátok az Intel C++-t?
    Vélemények róla?
    Mutasd a teljes hozzászólást!
  • thx. 99 usd nem sok erte, ha bevalik.
    Mutasd a teljes hozzászólást!
  • Próbáld ki a Visual Assist X-et refaktoráláshoz.
    Mutasd a teljes hozzászólást!
  • most mar a c++ elmeny is hasonlo a c#-hoz, csak a refaktoralas hianyzik. arrameg majd tuti lesznek addonok, ahogy eddig is. ami nagyon teccik, h a header es a cpp file is lenyithato, es latod felsorolva a deklaraciokat. mint anno a cpp builder, csak megjobb. meg azt kene valagogy megoldani, hogy a header es a fprras egy egyutt szerkesztodo egyseget kepezzen, hogy ne kelljen mindent 2x megrajzolni.

    edit: meg annyi hogy a boost jelenlegi verzioja nincs ra flekeszitve, de van a githubon egy meg nem mergelt changeset, ami hozzaadja a supportot.
    Mutasd a teljes hozzászólást!
  • Igen?
    Le is szedem...

    Hosszabb C# időszak után most megint C++ozok, ami azért elég nagy visszalépés az IDE szempontjából.
    Mutasd a teljes hozzászólást!
  • Igazság szerint én Linuxon dolgozom vim+clang_complete segítségével, illetve Qt Creatorral, de kösz :)
    Mutasd a teljes hozzászólást!
  • Nézd meg a VC 11 RC-t. 'va jó az IntelliSense, és végre (újra) van C++/CLI-hez is.

    Sz'al az MS télleg komolyan gondolta ezt a natív reneszánszt. :)
    Mutasd a teljes hozzászólást!
  • Mindenesetre: MinGW Distro - nuwen.net
    Ez a legteljesebb C++11 támogatást adó fordító jelenleg Windowsra. Persze így meg nincs VS, de igazából az intellisense C++-hoz amúgy sem túl jó benne, szóval mehet akár a Qt Creator is, abban jobb :)
    Mutasd a teljes hozzászólást!
  • Folytassuk itt. Ne legyen offtopic.
    Mutasd a teljes hozzászólást!
  • Mert szornyen korulmenyes, amig a proglamod eljut a vasig (kicsit hosszu lesz az explanation, amiert bocsesz :D):
    __kernel void f( __global uint *dst, __constant uint *cb) { dst[get_global_id(0)] = get_global_id(0) ^ cb[0];}

    Ebbol eloszor keszul egy LLVM_IR. Ez egy opensource fordito optimizer, ami kulonfele CPU style optimizaciok elvegzese utan egy tömör bytekódot allit elo.

    Az LLVM_IR alapján elkeszul az AMD_IL köztes nyelvre leforditott változat. (Hiába van ott a nevében, hogy ez köztes nyelv ez mar egy jo ideje hardvertol fuggoen kulonbozik) Erre az AMD_IL-re fordulnak tobbek kozott az opengl, directx, directcompute shaderek is.

    A fenti openCl kernel igy nez ki IL-ben:
    - felul szamitja ki a workitem indexeket
    - a call 1027-ben alul van a lenyeg
    mdef(16383)_out(1)_in(2) ;kotelezoen benne kell legyen, valami ocl okorseg miatt mov r0, in0 mov r1, in1 div_zeroop(infinity) r0.x___, r0.x, r1.x mov out0, r0 mend il_cs_2_0 dcl_cb cb0[15] ; Constant buffer that holds ABI data dcl_literal l0, 0x00000004, 0x00000001, 0x00000002, 0x00000003 dcl_literal l1, 0x00FFFFFF, 0xFFFFFFFF, 0xFFFFFFFE, 0xFFFFFFFD dcl_literal l2, 0x0000FFFF, 0xFFFFFFFE, 0x000000FF, 0xFFFFFFFC dcl_literal l3, 0x00000018, 0x00000010, 0x00000008, 0xFFFFFFFF dcl_literal l4, 0xFFFFFF00, 0xFFFF0000, 0xFF00FFFF, 0xFFFF00FF dcl_literal l5, 0x00000000, 0x00000004, 0x00000008, 0x0000000C dcl_literal l6, 0x00000020, 0x00000020, 0x00000020, 0x00000020 dcl_literal l7, 0x00000018, 0x0000001F, 0x00000010, 0x0000001F dcl_literal l8, 0x80000000, 0x80000000, 0x80000000, 0x80000000 call 1024;$ endmain func 1024 ; __OpenCL_F_kernel mov r1013, cb0[8].x mov r1019, l1.0000 dcl_max_thread_per_group 64 dcl_typeless_uav_id(10)_stride(4)_length(4)_access(read_write) mov r0.__z_, vThreadGrpIdFlat0.x mov r1022.xyz0, vTidInGrp0.xyz mov r1023.xyz0, vThreadGrpId0.xyz dcl_literal l9, 256, 1, 1, 0xFFFFFFFF umin r1023.xyz0, r1023.xyzz, l9.w ;obsolete umin r1021.xyz0, cb0[1].xyzz, l2.x ;fetch GroupSize imad r1021.xyz0, r1023.xyzz, r1021.xyzz, r1022.xyzz ;vThreadGrpId0*GroupSize+vTidInGrp0 umin r1024.xyz0, cb0[6].xyzz, l9.w ;obsolete iadd r1021.xyz0, r1021.xyz0, r1024.xyz0 ;shift? global_id() umin r1024.xyz0, cb0[7].xyzz, l9.w ;obsolete iadd r1023.xyz0, r1023.xyz0, r1024.xyz0 ;shift? group_id() mov r1023.___w, r0.z ishl r1023.___w, r1023.w, l0.z ;vThreadGrpIdFlat0.x shl 3 mov r1018.x___, l0.0000 udiv r1024.xyz_, r1021.xyzz, cb0[10].xyzz imad r1025.xyz0, r1023.xyzz, cb0[10].xyzz, r1022.xyzz dcl_literal l14, 0x00000000, 0x00000000, 0x00000000, 0x00000000; f32:i32 0 dcl_literal l11, 0x00000002, 0x00000002, 0x00000002, 0x00000002; f32:i32 2 dcl_literal l12, 0x00000003, 0x00000003, 0x00000003, 0x00000003; f32:i32 3 dcl_literal l15, 0x00010000, 0x00010000, 0x00010000, 0x00010000; f32:i32 65536 dcl_literal l13, 0x00000000, 0xFFFFFFFF, 0xFFFFFFFE, 0xFFFFFFFD; f128:i128 1844674406941458432018446744065119617022 dcl_cb cb1[2] ; Kernel arg setup: dst mov r1.x, cb1[0].x ;uav base address ; Kernel arg setup: cb mov r1.y, l5.x dcl_cb cb2[4096] call 1027 ; F ret ; commenthez nem nyulni, dolgozik belole!!!!!!! endfunc ; __OpenCL_F_kernel ;ARGSTART:__OpenCL_F_kernel ;version:3:1:104 ;device:tahiti ;uniqueid:1024 ;memory:uavprivate:0 ;memory:hwregion:0 ;memory:hwlocal:0 ;pointer:dst:u32:1:1:0:uav:10:4:RW:0:0 ;pointer:cb:u32:1:1:16:hc:2:4:RO:0:0 ;function:1:1027 ;privateid:8 ;reflection:0:uint* ;reflection:1:uint* ;ARGEND:__OpenCL_F_kernel func 1027 ; F ; @__OpenCL_F_kernel ; BB#0: ; %entry ; kernel opencl source code: dst[get_global_id(0)] = get_global_id(0) ^ cb[0] ; r1.x = uav base offset ; r1.y = cb2 base offset, seems like always zero ; r1021.x = get_global_id(0) ; ----------------------------------------- ; fetch constantBuffer[0] ; mov r1010.x___, r1.y ;get index (0) ; ishr r1007.x___, r1010.x, l11 ;shr 2 converts to dword offset ; iand r1008.x___, r1007.x, l12 ;and 3 get xyzw index ; ishr r1007.x___, r1007.x, l11 ;srh 2 converts to dqword offset ; iadd r1008, r1008.x, l13 ;adds (0,-1,-2,-3) shifts to 0 for ieq ; ieq r1008, r1008, l14 ;create xyzw selector mask ; mov r1011, cb2[r1007.x] ;loads the dqword ; cmov_logical r1011.x___, r1008.y, r1011.y, r1011.x ;choose the dword form the dqword ; cmov_logical r1011.x___, r1008.z, r1011.z, r1011.x ;choose the dword form the dqword ; cmov_logical r1011.x___, r1008.w, r1011.w, r1011.x ;choose the dword form the dqword ; calculate global_id(0) xor cb[0] ; mov r65.x___, r1011.x ;cb[0] ; mov r66, r1021.xyz0 ;global_id ; mov r65.__z_, r66.00x0 ;redundant move ; ixor r65.x___, r65.z, r65.x ;r65.x: global_id(0) xor cb[0] ; calculate address (byte offset pointing to a dword array) ; mov r65.___w, l11 ;redundant move ; ishl r65.__z_, r65.z, r65.w ;global_id(0)shl 2 ;adderss ; iadd r65._y__, r1.x, r65.z ;r65.y: add uav base offset(r1.x) ; UAV write ; mov r1011, r65.x ; value ; mov r1010.x___, r65.y ; address ; uav_raw_store_id(10) mem0.x___, r1010.x, r1011.x ret_dyn ret endfunc ; F ;ARGSTART:F ;uniqueid:1027 ;ARGEND:F end

    A fenti F funkciot human brain technologia segitsegevel le lehet redukalni ennyire:
    dcl_literal l99, 4, 0, 0, 0 ixor r1011.x, r1021.x, cb2[0].x //global_id(0)^cb[0] umad r1010.x, r1021.x, l99.x, r1.x //global_id(0)*4+dst.base address calculation uav_raw_store_id(10) mem0.x___, r1010.x, r1011.x //write

    Ez utan a szint utan jon az CAL compiler, ami leforditja hardver specifikus mikrokodra a fenti nyomokban hardver specifikus mikrogodot. Ez a fordito amugy fantasztikusan jol optimizal es minden konstanssal kapcsolatos szamitast kiszed.

    A hw spec. forras a kovetkezo:
    (Becsult orajelek baloldalon)
    0 s_load_dwordx4 s[12:15], s[2:3], 0x00 1 s_load_dwordx4 s[16:19], s[2:3], 0x04 2 s_load_dwordx4 s[0:3], s[2:3], 0x08 3 s_waitcnt lgkmcnt(0) 4 s_buffer_load_dword s9, s[12:15], 0x04 5 s_buffer_load_dword s10, s[12:15], 0x18 6 s_buffer_load_dword s11, s[16:19], 0x00 7 s_buffer_load_dword s0, s[0:3], 0x00 //fetch cb[0] 8 s_waitcnt lgkmcnt(0) 9 s_min_u32 s1, s9, 0x0000ffff 9 v_mov_b32 v1, s1 10 v_mul_i32_i24 v1, s8, v1 11 v_add_i32 v0, vcc, v0, v1 12 v_add_i32 v0, vcc, s10, v0 //v0-ban van vegre a global_id(0) 13 v_xor_b32 v1, s0, v0 //v1=global_id(0)^cb[0] 14 v_lshlrev_b32 v0, 2, v0 //addr=global_id(0)*4 15 v_add_i32 v0, vcc, s11, v0 //addr+=dst.base 15 tbuffer_store_format_x v1, v0, s[4:7], 0 offen format:[BUF_DATA_FORMAT_32,BUF_NUM_FORMAT_FLOAT] //write 16 s_endpgm

    Az jol latszik, hogy a szamitasi muveletbol a konstansokat szepen eltavolitja: pl. a cb[0] eloallitasa nagyon koltseges, mert 16byte aligned readdal kenytelen kiolvasni a 4 byteos [0] indexu dword-ot, ehhez kell plusz aritmetika, meg 3 conditional move. Mivel konstans az index [0] ezert ezt szepen megoldotta.
    Tehat amit lehet, szepen megcsinal, de nagyon keves informacioja van a feladatrol, emiatt kenytelen általánossagokat forditani es a legrosszabbra felkeszulni.

    Ami itt nem tul szep:
    - Minek kell 3 buffer objektummal szenvedni, amikor boven eleg az adott kernelnek 2 buffer objektum (global, const).
    - Tovabba, latja a fordito azt is, hogy untig eleg egy get_global_id(0) a kernelnek, de ez csakazertis kiszamolja a global/local/group/workitem id-ket x,y,z koordinatakra, sőt mindegyiket range-checkeli is :D A fenti ket dologgal tokol 14 utasitason keresztul. (ennek nagy reszet szepen kioptolja azert az amd_il -> isa fordito)
    - az s_load_ kepes arra, hogy egyetlen utasitassal beolvasson 16 dwordot is, de mivel a felsobb layerek (amd_il, opencl) nem sokat tudnak errol, ezert kenytelen dwordonkent olvasni az adatokat, holott ha szepen ossze lenne keszitve minden kernel parameter, akkor eleg lenne egy s_load_dwordx16 is.
    - nem hasznalhat 1 utasitasos mul_add-ot a base+offset*4 cim kiszamitashoz, mert nem tudja hogy az offset 24-biten elfer-e vagy nem (ilyen infot nem kaphat felulrol), pedig ha errol a kriteriumrol tudna, akkor ezt a cimszamitast meg tudná oldani az SP-float MAD tranzisztoraival.

    Hogy is lehetne ezt szepen megcsinalni CAL api alól?
    //{ dst[get_global_id(0)] = get_global_id(0) ^ cb[0];} 0 s_buffer_load_dword s0, s[12:15], 0x00 ;fetch cb[0] 0 v_mad_u24_u32 v0, s16, 64, v0 ;v0-ban a get_global_id(0), 64 waveitems/grp 1 s_waitcnt lgkmcnt(0) ;bevárjuk a cb[0] beolvasasat 1 v_xor_b32 v1, v0, s0 ;v1: get_global_id(0)^cb[0] 2 v_lshlrev_b32 v0, 2, v0 ;v0-ban a dword offset (base mindig 0 a CAL-ban, emiatt még muladd sem kell ide) 2 tbuffer_store_format_x v1, v0, s[4:7], 0 offen format:[BUF_DATA_FORMAT_32,BUF_NUM_FORMAT_FLOAT] ;vegul a write

    Az teljesen vilagos, hogy az OCL mennyire produktiv, csak hát nem igazan lehet benne ugy programozni, hogy kozben a 'lelki szemeiddel' látod is, hogy mit csinal a vason es tudod, hogy azt optimalisan is csinalja.

    No offense, de ez az amd-opencl ez így a gpu programozas .net-je azzal a 'nehezitessel', hogy itt nincs unsafe mode.

    Meg egy konkret dolog, de aztan most mar tenyleg abbahagyom haha:
    Tehat az uj radeonok igencsak mashogy mukodnek, mint az eddig parhuzamos architekturak, ehhez uj fajta gondolkodasmod kell:

    A konstansok elokesziteseert most mar egy scalar alu felel, ami 1 muveletet vegrehajt, mialatt a vektor alu 64 parhuzamos muveletet.
    De a skalar alu nem csak a konstansok elokaparasara kepes, hanem tud szamitasokat is elvegezni.
    Ez pedig jol hasznalva szamitasokat vehet át a vektor alu-tol: Ha sikerul ugy csoportositani az egy workgroupban levo workitemeket, hogy az egyes workgroupokon belul sok olyan megegyezo eredmenyu szamitas legyen, akkor ezeket a feladatokat at lehet passzolni a scalar alu-nak.
    Tovabb bonyolitja a helyzetet, hogy a scalar alu 64biten is tud szamolni :D Tehat azon a sok biten mar lehet 'kamu' simd-zni is. Peldaul ha logikai muveletel vesz át a vektor alu-tol, akkor azt dupla olyan gyorsan tudja daralni, csak egymas melle kell parositani a 32bit reszadatokat.

    Szoval eddig ugy voltam vele, hogy a VLIW tipusu gpu-khoz jo lesz nekem a ams_il koztesnyelv is. De az uj 7ezres radeonnal mar a mikrokod is van olyan egyszeru, mint az intel asm ( csak van benne scalar/vektor alu csavar).
    Szoval en a jovoben arra fogok raalni, hogy a magas szintu ocl api-bol paszirozzam a gpu-ba a legalacsonyabb szintu asm-ot haha.

    Ez persze csak a radeon oldalra vonatkozik, de a geforcenal is hasonlo lehet a helyzet, hogy inkabb a cuda, mert azt tenyleg ahhoz a vashoz talaltak ki. De az ocl nyelv meg egy kepzeletbeli vashoz lett kitalalva akkoriban, amikor meg olyan hardver nem is volt, mint most.

    De hogy pozitivum is legyen azert :D Az opencl api vegre a catalyst dll-jein beluli accessviolak helyett konkret hibauzeneket is tud mondani. Es nem nagyon tudok OS-t fagyasztani vele fejlesztes kozben :D Az alatta levo frontend-el (cal) meg mar nem foglalkoznak, de vegulis itt is mindent meg tudok oldani, amit ott...

    Na end of regény
    Mutasd a teljes hozzászólást!
  • ineropol helyett pinvoke-olt akartam irni. az megtetubb.

    (bocs az egyedi helyesirasert, telefonrol kicsit nehez a magyar abc-t es a szoalkoto szabalyokat artikulalni)
    Mutasd a teljes hozzászólást!
  • mi a bajod az opencl-lel? en ennek koszonhetoen tanultam meg a gpgpu shaderek irasat. mikor belekezdem, nem mertem volna gondolni, hogy ez ennyire egyszeru is lehet. bar azota kijott a c++ amp, ami meg sokkal egyszerubb es elegansabb, bar egyelore eleg lassucska szegeny.

    a c++-t is az ocl miatt vettem elo, mert a cloo, vagyis a .nethez rajzolt ocl binder interopol, vagyis hot tetu felallasbol indit. az ocl c++ bindere ellenben nagyon elegans megoldas, pont annyit ad hozza a c apihoz, amennyit kell ahhoz, hogy konnyu legyen hasznalni, de a sebesseg is megmaradjon.

    aztan ha mar c++, felraktam a boostot, az megtetszett, es igy jott a tobbi. ez a facecebook lib elvileg a boost es az stl bizonyos reszeit adja, csak sokkal inkabb a sebessegre kihegyezve, mint az eleganciara, ezert vagyok ra nagyon kivancsi. fog kelleni.
    Mutasd a teljes hozzászólást!
  • Lol, te lefelé nézel c++ irányba, en meg most kenytelen vagyok egy szinttel feljebb OCL-re átállni, mert az alatta levo szint api-ját szisztematikusan nyírja kifele az amd.
    Mutasd a teljes hozzászólást!
  • indid. pont azer kezdtem el a hobbi projektemben komolyabban c++-ozni az utobbi evben, mert szvsz ez a win8 ize akkora fiasko lesz, h viszi magaval az egesz ceget a csod, es a .net helyett ezerszer inkabb cememeg min java. szerencsere a c++ amp nyilt szabvany.
    Mutasd a teljes hozzászólást!
  • Igen, az frankó, ellenben azt nagyon rosszul viselem, hogy csak metro fejlesztéshez adják ingyen a fordítót (ami C++/CX ráadásul). Még ha a Microsoft üzleti szemszögéből nézzük, akkor sem vagyok biztos benne, hogy ez bölcs/profitábilis stratégia.
    Mutasd a teljes hozzászólást!
  • Hátja. Igaz. Ez még elég hurka.

    Ellenben az MS 11-es fordítójához adott PPL, C++ AMP na meg a végre megcsinált auto vektorizáció olyan musthave, ami miatt a fenti lista felett én simán szemet tudok húnyni.

    (Halkan hozzátéve azért, hogy az RC-s C++ AMP adatmozgatási sebessége még mindig 40%-kal lassabb mint az OpenCL, gyakorlatilag 0%-os sebességoptimalizáció történt a béta óta, pedig nagyon ígérgették a fórumon, hogy ez már nagyon jó lesz. Talán végleges.)
    Mutasd a teljes hozzászólást!
  • Én inkább úgy fogalmaznék, hogy remélhetőleg az MS összekapja magát a C++ fordítója terén. Ez elég szabványos C++ kód, nem a library-t kellene megváltoztatni, hanem a fordítót képessé tenni a lefordítására.
    Mutasd a teljes hozzászólást!
  • pl a boost optionalba, variantba nagyon kene
    Mutasd a teljes hozzászólást!
  • Úgy emlékszem ez egy nagyon jó cikk!
    érthető
    Mutasd a teljes hozzászólást!
  • Köszönöm, így már tisztább a kép.
    Mutasd a teljes hozzászólást!
  • "c++ move semantics" keress rá, hasznos dolog.
    Mutasd a teljes hozzászólást!
  • Bocs, de mi az a "move kompatibilis"? Erről még nem hallottam. Google hirtelen nem segít, trillió PS3 Move kontrolleres találatból nehéz.
    Mutasd a teljes hozzászólást!
  • cool. remelhetoleg lesz belole vc 11-en es x64-en fordulo valtozat is. plusz nem artana ha ez mar 100%-ban move kompatibilis lenne, nem ugy mint nehany regebbi boost class.
    Mutasd a teljes hozzászólást!
  • Ami még fejlesztői szempontból fontos (a cikkből sajnos kimaradt), hogy mindezt Apache Licenc 2.0 alatt tette meg, tehát a kód szabadon felhasználható bármilyen projektben (el is adható).
    Mutasd a teljes hozzászólást!
Tetszett amit olvastál? Szeretnél a jövőben is értesülni a hasonló érdekességekről?
abcd