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));
}
#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);
}
#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);
}
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.
Care About Your Craft: Why spend your life developing software unless you care about doing it well? - The Pragmatic Programmer