File size: 2,436 Bytes
e873e70 | 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 | // Can the iGPU and CPU run relu at the same time? They share DDR5. This runs both
// concurrently on separate buffers and compares combined BW to each alone.
#include <sycl/sycl.hpp>
#include <omp.h>
#include <immintrin.h>
#include <cstdio>
#include <cstdlib>
#include <vector>
#include <chrono>
#include <algorithm>
using namespace sycl;
static void cpu_relu(float* __restrict o, const float* __restrict in, size_t n){
const __m256 z=_mm256_setzero_ps();
#pragma omp parallel for schedule(static)
for(size_t j=0;j<n;j+=8) _mm256_stream_ps(o+j,_mm256_max_ps(_mm256_loadu_ps(in+j),z));
_mm_sfence();
}
template<class F> static double best_of(F f,int reps=8){
for(int w=0;w<2;w++) f();
double best=1e30;
for(int r=0;r<reps;r++){ auto a=std::chrono::high_resolution_clock::now(); f();
auto b=std::chrono::high_resolution_clock::now();
best=std::min(best,std::chrono::duration<double>(b-a).count()); }
return best;
}
int main(){
queue q{gpu_selector_v};
printf("iGPU: %s + CPU: %d OpenMP threads\n",
q.get_device().get_info<info::device::name>().c_str(), omp_get_max_threads());
const size_t ng=64ull*1024*1024, nc=64ull*1024*1024;
float* gin =malloc_device<float>(ng,q); float* gout=malloc_device<float>(ng,q);
std::vector<float> hg(ng); for(size_t i=0;i<ng;++i) hg[i]=(i&1)?-1.f:1.f;
q.memcpy(gin,hg.data(),ng*4).wait();
float* cin =(float*)aligned_alloc(64,nc*4); float* cout=(float*)aligned_alloc(64,nc*4);
for(size_t i=0;i<nc;++i) cin[i]=(i&1)?-1.f:1.f;
auto gpu_submit=[&](){ q.parallel_for(range<1>(ng),[=](id<1> id){ size_t k=id[0]; float x=gin[k]; gout[k]=x>0.f?x:0.f; }); };
const double gbg=2.0*ng*4/1e9, gbc=2.0*nc*4/1e9;
double tg = best_of([&](){ gpu_submit(); q.wait(); }); // iGPU alone
double tc = best_of([&](){ cpu_relu(cout,cin,nc); }); // CPU alone
double tb = best_of([&](){ gpu_submit(); cpu_relu(cout,cin,nc); q.wait(); }); // BOTH at once
double bw_g=gbg/tg, bw_c=gbc/tc, bw_both=(gbg+gbc)/tb;
printf("\niGPU alone : %5.1f GB/s\n", bw_g);
printf("CPU alone : %5.1f GB/s\n", bw_c);
printf("BOTH concurrently : %5.1f GB/s combined\n", bw_both);
printf(" if independent : %5.1f GB/s (= sum)\n", bw_g+bw_c);
printf(" scaling vs sum : %.0f%% (100%% = no contention, ~50%% = fully memory-bound)\n",
100.0*bw_both/(bw_g+bw_c));
free(gin,q); free(gout,q); free(cin); free(cout);
return 0;
}
|