search

UoB-HPC / BabelStream

STREAM, for lots of devices written in many programming models
Other
323 stars 110 forks source link

Fixed sycl kernel linking for computecpp #87

Closed tom91136 closed 3 years ago

tom91136 commented 3 years ago

So it turns out there was an oversight for ComputeCpp in https://github.com/UoB-HPC/BabelStream/commit/f0403d2b092115a6b0146cd030ff4a006520d3a9 where the sycl kernel isn't actually included in the compiled binary. The effect of this is quite surprising: when ran on the host device, everything works, but if we choose any other SYCL device, it throws a sycl::invalid_object_error exception on the first submit call. I suspect the original commit was accidentally tested with the host device.

ComputeCpp's official Makefile template now uses the compute++ compiler to compile everything so we also default to that. I've looked into whether we can reuse the original dependency chain but I gave up and made it conditional instead; dissecting their CMake plugin, it shows further .bc and .sycl file dependencies which makes the process quite different from the others.

tom91136 commented 3 years ago

I've cleaned up the targets and variables now with all SYCL compilers getting the source list (just two files) so no more special case required for ComputeCpp.

Everything has been tested locally with a proper sycl device and the numbers are within ballpark.