[SYCL] Specialness of SYCL vec types and performance of Eigen in device code

[AnonymousRoboticist]

I am currently investigating using the Eigen library in the device code (not Eigen with SYCL backend) as my usecase involves a lot of small vector, matrix, quaternion, etc computations in parallel. I have found that I can get Eigen to run in device code using

// Avoid automatic memory allocations
#define EIGEN_RUNTIME_NO_MALLOC
#define EIGEN_NO_MALLOC
// Disable sse2/etc vectorization SIMDs for spirv gen
#define EIGEN_DONT_VECTORIZE
#include <Eigen/Dense>

What I was wondering was if there's anything "special" to sycl vec types when it comes to performance. Are they simply types that integrate well and are specified with the rest of SYCL, or do they have special behaviour or heuristics programmed in the compiler for optimiser. Will replacing the use of sycl::float4 with Eigen::Vector4f result in significant performance differences?

[AlexeySachkov]

Hi @AnonymousRoboticist,

As SYCL spec says, vec class is intended to map to some efficient representation on a target HW:

SYCL provides a cross-platform class template that works efficiently on SYCL devices as well as in host C++ code. This type allows sharing of vectors between the host and its SYCL devices. The vector supports member functions that allow construction of a new vector from a swizzled set of component elements.

vec<typename DataT, int NumElements> is a vector type that compiles down to a SYCL backend built-in vector types on SYCL devices, where possible, and provides compatible support on the host or when it is not possible.

I suppose that for Eigen (or any other project) there is always a choice: either have custom efficient implementations for all target HW or use a standard functionality and rely on SYCL implementation to make the most out of it. Also tagging @rolandschulz here.

Will replacing the use of sycl::float4 with Eigen::Vector4f result in significant performance differences?

It is hard to say for me, but it could be the case depending on how Eigen::Vector4f is implemented. sycl::float4 bears the semantic of being a vector data type, so each BE can easily map it to some specific HW instructions. With a custom user-defined data structure compiler might not be able to figure out the intent of that structure and therefore might miss some optimization opportunities, I suppose.