Intels Xe-grafikkort skymtas – instegskort med 3 GB grafikminne

Permalänk
Datavetare

Instegskort i alla ära... Det verkar ändå som Intels primära fokus runt Xe kretsar kring GPGPU, inte spelprestanda.

Nvidias VD gick redan i höstas ut och dissade Intels plattform som riktar sig mot programmering av SIMD på CPU, GPUer och FPGA:er. Svårt att säga så mycket i det läget då det inte fanns så mycket mer än PowerPoint.

Men sedan början av året har Intel publicerat beta-releaser av deras oneAPI. Intel har definitivt uppförsbacke framför sig givet hur Nvidia totalt dominerar GPGPU med sitt CUDA-ramverk, men Jensen får nog fundera lite för till skillnad från OpenCL som var allt för komplicerat att använda jämfört med CUDA så ser oneAPI riktigt intressant ut!

OpenCL hade flera designmissar, alla är åtgärdade i oneAPI (som till stor del bygger på en öppen standard från Khronos som kallas SYCL). Men att göra något som är lika bra som CUDA tror jag inte räcker, AMD försöker med precis detta med deras HIP som ligger väldigt nära CUDA.

De som designade SYCL, vilket inte var Intel utan ett företag som heter CodePlay, tog verkligen ett steg tillbaka och funderade på vilka tillkortakommanden som fanns i CUDA/OpenCL. Ovanpå sin expertis i heterogena system var CodePlay insiktsfulla att också plocka in experter på C++ och kompilatorteknik i designen. Att programmera parallella system är svårt nog som det är, att hantera multipla program (program som körs i CUDA, SYCL och HIP kallas "kernels") som i sig är massivt parallell och sedan kör parallellt med andra "kernels" är än mer komplicerat.

Just här finns fördelar i SYCL/oneAPI saknas i CUDA/HIP. Det är möjligt att göra motsvarande i de senare, men det händer inte av sig själv utan en människa måste explicit inse hur olika kernels använder data och därmed hur de beror av varandra. Gör man fel där får man en av de svåraste buggar som finns att leta efter: data-race. SYCL/oneAPI har en design som gör att databeroenden kan redas ut redan i kompilatorsteget och kompilatorn, inte en människa, räknar ut hur saker ska köras och synkroniseras i relation till varandra.

Att addera alla element i potentiellt stora vektorer är lite "Hello world" för GPGPU. Att göra det i OpenCL innehåller så mycket boilerplate att man blir rent matt, men redan ett sådant simpelt exempel visar några av fördelarna med SYCL (utöver att exakt samma kod kan köras på GPU, CPU och FPGA, något som inte är möjligt med CUDA/HIP).

#include <cstdio> #include <cmath> #include <vector> #include <CL/sycl.hpp> namespace sycl = cl::sycl; int main(int argc, char * argv[]) { // Number of elements in vector size_t n = 20; // Define the 1D vectors as standard C++ vectors std::vector<float> a_h(n); std::vector<float> b_h(n); std::vector<float> c_h(n); // Fill the source vectors with some well defined data for (auto i = 0; i < n; i++) { a_h[i] = sin(i) * sin(i); b_h[i] = cos(i) * cos(i); } // Define device buffers on top of C++ vectors sycl::buffer<float, 1> a_buf(a_h.data(), a_h.size()); sycl::buffer<float, 1> b_buf(b_h.data(), b_h.size()); sycl::buffer<float, 1> c_buf(c_h.data(), c_h.size()); // Use the "default" SYCL device sycl::queue device_q; // Run kernel on device, kernel can be defined as C++11 lambda! device_q.submit([&](sycl::handler& h) { // Reqeust access to buffers from the device auto A = a_buf.get_access<sycl::access::mode::read>(h); auto B = b_buf.get_access<sycl::access::mode::read>(h); auto C = c_buf.get_access<sycl::access::mode::write>(h); h.parallel_for<class vec_add>( sycl::range<1>(n), [=] (sycl::id<1> idx) { C[idx] = A[idx] + B[idx]; } ); }); // Request access to buffer c_buf from CPU, SYCL compiler will figure out // data-dependency between vec_add kernel and printing the result! c_buf.get_access<sycl::access::mode::read>(); // Sum all elements in c_h and print the result std::printf("Sum: %f\n", std::accumulate(begin(c_h), end(c_h), 0.0)); }

oneAPI

#include <stdio.h> #include <stdlib.h> #include <math.h> // CUDA kernel. Each thread takes care of one element of c __global__ void vecAdd(float *A, float *B, float *C, int N) { // Get our global thread ID int idx = threadIdx.x; // Make sure we do not go out of bounds if (idx < N) { C[idx] = A[idx] + B[idx]; } } int main( int argc, char* argv[] ) { // Number of elements in vector int n = 20; // Host input vectors float *h_a; float *h_b; //Host output vector float *h_c; // Device input vectors float *d_a; float *d_b; //Device output vector float *d_c; // Size, in bytes, of each vector size_t bytes = n * sizeof(float); // Allocate memory for each vector on host h_a = (float*)malloc(bytes); h_b = (float*)malloc(bytes); h_c = (float*)malloc(bytes); // Allocate memory for each vector on GPU cudaMalloc(&d_a, bytes); cudaMalloc(&d_b, bytes); cudaMalloc(&d_c, bytes); // Fill the source vectors with some well defined data for (size_t i = 0; i < n; i++) { h_a[i] = sin(i) * sin(i); h_b[i] = cos(i) * cos(i); } // Copy host vectors to device cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice); cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice); int blockSize, gridSize; // Number of threads in each thread block blockSize = 1024; // Number of thread blocks in grid gridSize = (int)ceil((float)n/blockSize); // Execute the kernel, no CUDA-stream defined, implicitly using stream 0 vecAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n); // Copy array back to host cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost); // Sum up vector c and print result divided by n, this should equal 1 within error float sum = 0; for(size_t i = 0; i < n; i++) { sum += h_c[i]; } printf("Sum: %.2f\n", sum); // Release device memory cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); // Release host memory free(h_a); free(h_b); free(h_c); }

CUDA

#include <stdio.h> #include <stdlib.h> #include <math.h> // HIP kernel. Each thread takes care of one element of c __global__ void vecAdd(const float *A, const float *B, float *C, int N) { // Get our global thread ID int idx = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; // Make sure we do not go out of bounds if (idx < N) { C[idx] = A[idx] + B[idx]; } } int main( int argc, char* argv[] ) { // Number of elements in vector int n = 20; // Host input vectors float *h_a; float *h_b; //Host output vector float *h_c; // Device input vectors float *d_a; float *d_b; //Device output vector float *d_c; // Size, in bytes, of each vector size_t bytes = n * sizeof(float); // Allocate memory for each vector on host h_a = (float*)malloc(bytes); h_b = (float*)malloc(bytes); h_c = (float*)malloc(bytes); // Allocate memory for each vector on GPU hipMalloc((void **)&d_a, bytes); hipMalloc((void **)&d_b, bytes); hipMalloc((void **)&d_c, bytes); // Fill the source vectors with some well defined data for (size_t i = 0; i < n; i++) { h_a[i] = sin(i) * sin(i); h_b[i] = cos(i) * cos(i); } // Copy host vectors to device hipMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice); hipMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice); int blockSize, gridSize; // Number of threads in each thread block blockSize = 1024; // Number of thread blocks in grid gridSize = (int)ceil((float)n/blockSize); // Execute the kernel, no HIP-stream defined, implicitly using stream 0 hipLaunchKernelGGL(vecAdd, dim3(gridSize), dim3(blockSize), 0, 0, d_a, d_b, d_c, n); // Copy array back to host hipMemcpy(h_c, d_c, bytes, hipMemcpyDeviceToHost); // Sum up vector c and print result divided by n, this should equal 1 within error float sum = 0; for(size_t i = 0; i < n; i++) { sum += h_c[i]; } printf("Sum: %.2f\n", sum/n); // Release device memory hipFree(d_a); hipFree(d_b); hipFree(d_c); // Release host memory free(h_a); free(h_b); free(h_c); }

HIP, notera HIP kräver ROCm drivaren som bara finns på Linux och det saknas stöd för rDNA än så länge

oneAPI varianten är kortare, ändå innehåller den möjligheter som saknas i CUDA/HIP versionen. T.ex. möjlighet att köra flera "kernels" parallelt (krävs att man explicit definierar CUDA-streams för det) samt möjlighet att köra på CPU (t.ex. för enklare debugg och/eller för att man gör något som mer effektivt hanteras med SSE/AVX/AVX-512 jämfört med GPGPU).

Just att samma kod fungerar både på GPU och CPU ihop med automatisk beroendehantering gör det någorlunda enkelt att använda CPU och GPU parallellt. Det är tekniskt sett möjligt även med CUDA/HIP, men man får hantera synkroniseringen manuellt samt man måste skriva en version för GPU och en för CPU.

Framtiden lär visa om oneAPI är bra nog. Det räcker inte att ha bra HW, utan bra programvara är HW meningslös! Nvidia dominerar över AMD på GPGPU just då man haft långt bättre programvara, AMD har vid flera tillfällen haft övertag i rå HW-prestanda på GPGPU-sidan.

Visa signatur

Care About Your Craft: Why spend your life developing software unless you care about doing it well? - The Pragmatic Programmer

Permalänk
Datavetare
Skrivet av sKRUVARN:

Jag pratar om instickskort för stationära nu, alltså det som visas på bild i artikeln, och det handlade inte den länken om, står specifikt mobile och laptop i var och varannan mening. Och vi kan ju jämföra den designen med ett gt720 med 20w rtp

https://tpucdn.com/gpu-specs/images/c/2702-front.jpg

Edit. Just det där kortet är nog dock lite nerklockat för att komma undan med sådan liten kylare, men oavsett kolla på kortdesign på kort runt denna tdp, kortet som visas i artikeln är betydligt bitigare. Och som sagt, dessa kort är rätt utspelade vid det här laget, finns inte mycket poäng att släppa ett instickskort (eller ja, nån nytta finns såklart) som man ska spela på som inte är mycket snabbare än igpu.

Edit2.
Kanske kan droppa diskussionen, ser inte ut att bli något instickskort?

https://www.techpowerup.com/266920/intel-xe-dg1-silicon-not-m...

Att DG1-LP rimligen bara skulle riktas mot bärbara kändes som enda rimliga redan när man visade upp DG1LPDev på CES. Påpekade det absurda i den "analys" flera tech-tubers försökte sig på kring spelprestanda hos DG1LPDev, men var uppenbarligen många som verkade tro att det var kortet Intel skulle ge sig på AMD/Nvidia med...

Vad man sagt är ju att Tiger Lake ska få ungefär dubbla GPU-prestanda mot Ice Lake. Om "IPC" är ungefär densamma för gen11 och gen12 kan bara detta uppnås om man når 1,3-1,4 GHz. Möjligen kan man klocka en separat krets ännu lite högre då den inte behöver dela TDP på samma sätt med CPU, men om man nu nått 1,5 GHz är det ändå i en krets som rimligen primärt riktas mot MX250/MX350 och därmed borde få en TDP på 10-25 W då det är området Nvidias kretsar rör sig mellan (exakt TDP beror på hur slimmad dator man stoppar in kretsen i).

Vid 1,5 GHz och med 96 EU når man nästan (men bara nästan) den teoretiska prestandan hos GTX1650. Så fortfarande absolut instegsmodell, men tillräckligt för att Nvidia kanske får skrota MX350 innan den knappt lanserats (MX150/MX250 ~ GT1030, MX350 ~ GTX1050)

Så 25 W känns som ett tak. Givet att DG1LPDev kortet bara gick på 1,05 GHz borde den i alla fall inte dragit mer.

Har denna GT1030, den drar verkligen 30 W vid full last (har mätt detta och är också vad kortet rapporterar via nvidia-smi). Bara man har lite luftflöde så fixar den kylningen, men är det helt stiltje runt kortet verkar det överhetta (datorn kan hänga sig) under väldigt långvarig GPGPU-last (har aldrig sett det hänga sig vid grafiklast, men inte precis ett spelkort så inte testat så ofta...).

Ändå än rätt redig kylare! Och det är en sak att fixa vårat klimat med 20-25 grader inomhus, men man ska ju typisk klarar lite högre omgivande temperatur än så i verkliga konsumentprodukter.

Visa signatur

Care About Your Craft: Why spend your life developing software unless you care about doing it well? - The Pragmatic Programmer

Permalänk
Avstängd

Dom har designat ett Alibaba GTX 650🤦.

Visa signatur

Alldeles för svag för SFF 🥰

Permalänk
Medlem
Skrivet av Sando:

Definitivt. Intel behöver inte nå upp till RTX 2080 Ti-nivå för att ha en påverkan på priserna i toppsegmentet. Om de t.ex. kunde nå upp till RTX 2060-nivå för konkurrenskraftiga priser (säkert minusaffärer i början) så skulle det tvinga ner priserna i det segmentet. Om man plötsligt får RTX 2060-prestanda för 2000kr så blir Nvidia och AMD tvungna att sänka priserna i de segmenten, vilket gör att segmentet över måste sänkas för att hålla okej prisvärdhet. Sen blir det en kedjeeffekt genom hela marknaden.

Förhoppningsvis kommer vi snart igen till en tid då toppkorten kostar kanske 700 usd, och inte 1000 usd i teorin och 1200 usd i praktiken.

Låter underbart! Gogo Intel!

Visa signatur

SpelDator: | Asus ROG Strix XG438Q 4K 120hz| AMD Ryzen 5 5600X 3.7 GHz 35MB| B450 I AORUS PRO WIFI| MSI Radeon RX 6800 16GB | Corsair Force MP510 1920GB M.2 NVMe | Samsung EVO 860 2TB | Seagate Guardian BarraCuda 2.5 5TB| Corsair 32GB (2x16GB) DDR4 3600MHz CL18 Vengeance LPX Svart| Ghost S1 MK2| Corsair SF750 750W 80+ Platinum| Logitech MX Sound| WD My Passport V2 USB 3.0 4TB| Buffalo Mediastation Ultra-Thin Portable Blu-Ray writer| Logitech Logitech G604 Lightspeed

Permalänk
Medlem

Intel Xe verkar bli producerad på 5 nm.

Om man nu ska tro TSMC produktplan som publicerats i China Times.

Bild och information via IgorsLAB

Visa signatur

Engineer who prefer thinking out of the box and isn't fishing likes, fishing likes is like fishing proudness for those without ;-)
If U don't like it, bite the dust :D
--
I can Explain it to you, but I can't Understand it for you!

Permalänk

Hoppas på bra prestanda i både FP16, FP32 och självklart även FP64.

Visa signatur

Dator: EEE901 N270/ 1GB / 20GB SSD ... Kraftigt nedbantat/tweakat Win7 x86 for speeeEED!
Facebook användare? Hatar tidslinjen? gå med i denna FB-grupp:
Undo Timeline ...med lite tur får h*lvetet ett slut!