From 7318daf887c4f06fa62e59e29fa675e48ad168f9 Mon Sep 17 00:00:00 2001 From: Mehdi Goli Date: Fri, 25 Nov 2016 16:19:07 +0000 Subject: Fixing LLVM error on TensorMorphingSycl.h on GPU; fixing int64_t crash for tensor_broadcast_sycl on GPU; adding get_sycl_supported_devices() on syclDevice.h. --- unsupported/Eigen/CXX11/src/Tensor/TensorIntDiv.h | 8 ++++++++ 1 file changed, 8 insertions(+) (limited to 'unsupported/Eigen/CXX11/src/Tensor/TensorIntDiv.h') diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorIntDiv.h b/unsupported/Eigen/CXX11/src/Tensor/TensorIntDiv.h index ede3939c2..eea25ac33 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorIntDiv.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorIntDiv.h @@ -37,6 +37,8 @@ namespace { { #ifdef __CUDA_ARCH__ return __clz(val); +#elif defined(__SYCL_DEVICE_ONLY__) + return cl::sycl::clz(val); #elif EIGEN_COMP_MSVC unsigned long index; _BitScanReverse(&index, val); @@ -53,6 +55,8 @@ namespace { { #ifdef __CUDA_ARCH__ return __clzll(val); +#elif defined(__SYCL_DEVICE_ONLY__) + return cl::sycl::clz(val); #elif EIGEN_COMP_MSVC && EIGEN_ARCH_x86_64 unsigned long index; _BitScanReverse64(&index, val); @@ -88,6 +92,8 @@ namespace { EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE uint32_t muluh(const uint32_t a, const T b) { #if defined(__CUDA_ARCH__) return __umulhi(a, b); +#elif defined(__SYCL_DEVICE_ONLY__) + return cl::sycl::mul_hi(a, static_cast(b)); #else return (static_cast(a) * b) >> 32; #endif @@ -97,6 +103,8 @@ namespace { EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE uint64_t muluh(const uint64_t a, const T b) { #if defined(__CUDA_ARCH__) return __umul64hi(a, b); +#elif defined(__SYCL_DEVICE_ONLY__) + return cl::sycl::mul_hi(a, static_cast(b)); #elif defined(__SIZEOF_INT128__) __uint128_t v = static_cast<__uint128_t>(a) * static_cast<__uint128_t>(b); return static_cast(v >> 64); -- cgit v1.2.3