Kernels
relu / bench /sycl_igpu_modes.cpp
superexpai's picture
Optimized relu: cpu/cuda/xpu, 1.2-1.85x faster on RTX 4090, benchmarked vs upstream and torch.relu
e873e70 verified
Raw
History Blame Contribute Delete
2.22 kB
// 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 <sycl/sycl.hpp>
#include <cstdio>
#include <cstdint>
#include <vector>
#include <chrono>
#include <algorithm>
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<class T> T fillval(size_t i);
template<> float fillval<float>(size_t i) { return ((i&1)?-1.f:1.f)*float(i%97); }
template<> uint16_t fillval<uint16_t>(size_t i){ return (uint16_t)((i&1)?(0x8000u|(i%200)):(i%200)); }
template<> int8_t fillval<int8_t>(size_t i) { return (int8_t)((i&1)?-(int)(i%100):(int)(i%100)); }
template<typename T>
void run(sycl::queue& q, const char* name) {
const size_t n = 64ull*1024*1024;
T* in = malloc_device<T>(n, q);
T* out = malloc_device<T>(n, q);
std::vector<T> h(n), o(n);
for (size_t i=0;i<n;++i) h[i]=fillval<T>(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<n&&ok;++i) if(o[i]!=relu_dev(h[i])) ok=false;
double best=1e30;
for(int r=0;r<12;++r){ auto a=std::chrono::high_resolution_clock::now(); rn();
auto b=std::chrono::high_resolution_clock::now();
best=std::min(best,std::chrono::duration<double>(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<info::device::name>().c_str(),
q.get_device().get_info<info::device::max_compute_units>());
printf("extended XPU relu, all dtypes (64M):\n");
run<float>(q, "fp32");
run<uint16_t>(q, "fp16/bf16"); // identical 16-bit sign-bit op
run<int8_t>(q, "int8");
return 0;
}