// Extended XPU relu kernel (fp32/fp16/bf16/int8) run on the Intel iGPU via SYCL. // Same logic as relu_xpu/relu.cpp: max for fp32/int8, sign-bit zero for 16-bit. #include #include #include #include #include #include using namespace sycl; inline float relu_dev(float x) { return x > 0.f ? x : 0.f; } inline uint16_t relu_dev(uint16_t x) { return (x & 0x8000u) ? uint16_t(0) : x; } // fp16 & bf16 inline int8_t relu_dev(int8_t x) { return x > 0 ? x : int8_t(0); } template T fillval(size_t i); template<> float fillval(size_t i) { return ((i&1)?-1.f:1.f)*float(i%97); } template<> uint16_t fillval(size_t i){ return (uint16_t)((i&1)?(0x8000u|(i%200)):(i%200)); } template<> int8_t fillval(size_t i) { return (int8_t)((i&1)?-(int)(i%100):(int)(i%100)); } template void run(sycl::queue& q, const char* name) { const size_t n = 64ull*1024*1024; T* in = malloc_device(n, q); T* out = malloc_device(n, q); std::vector h(n), o(n); for (size_t i=0;i(i); q.memcpy(in, h.data(), n*sizeof(T)).wait(); auto rn=[&](){ q.parallel_for(range<1>(n), [=](id<1> idx){ size_t i=idx[0]; out[i]=relu_dev(in[i]); }).wait(); }; for(int w=0;w<5;++w) rn(); q.memcpy(o.data(), out, n*sizeof(T)).wait(); bool ok=true; for(size_t i=0;i(b-a).count()); } double gb=2.0*n*sizeof(T)/1e9; printf(" %-12s %dB %5.1f GB/s %6.1f Gel/s %s\n", name, (int)sizeof(T), gb/best, n/best/1e9, ok?"OK":"FAIL"); free(in,q); free(out,q); } int main(){ queue q{gpu_selector_v}; printf("iGPU: %s (%u EUs)\n", q.get_device().get_info().c_str(), q.get_device().get_info()); printf("extended XPU relu, all dtypes (64M):\n"); run(q, "fp32"); run(q, "fp16/bf16"); // identical 16-bit sign-bit op run(q, "int8"); return 0; }