From 06eb24cf4d7d54e56abfb37ea062a7cb0c887550 Mon Sep 17 00:00:00 2001
From: Gael Guennebaud <g.gael@free.fr>
Date: Fri, 13 Jul 2018 16:04:27 +0200
Subject: Introduce gpu_assert for assertion in device-code, and disable them
 with clang-cuda.

---
 unsupported/Eigen/CXX11/src/Tensor/TensorConvolutionSycl.h | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

(limited to 'unsupported/Eigen/CXX11/src/Tensor/TensorConvolutionSycl.h')

diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorConvolutionSycl.h b/unsupported/Eigen/CXX11/src/Tensor/TensorConvolutionSycl.h
index da88bcb3b..65403905a 100644
--- a/unsupported/Eigen/CXX11/src/Tensor/TensorConvolutionSycl.h
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorConvolutionSycl.h
@@ -352,7 +352,7 @@ struct TensorEvaluator<const TensorConvolutionOp<Indices, InputArgType, KernelAr
           size_t range_x, GRange_x, tileSize_x, range_y, GRange_y, tileSize_y;
           m_device.parallel_for_setup(numX, numP, tileSize_x,tileSize_y,range_x,range_y, GRange_x, GRange_y );
           const size_t shared_mem =(tileSize_x +kernel_size -1)*(tileSize_y);
-          assert(static_cast<unsigned long>(shared_mem) <= m_device.sharedMemPerBlock());
+          gpu_assert(static_cast<unsigned long>(shared_mem) <= m_device.sharedMemPerBlock());
           auto global_range=cl::sycl::range<2>(GRange_x, GRange_y);  // global range
           auto local_range=cl::sycl::range<2>(tileSize_x, tileSize_y);  // local range
           InputLocalAcc local_acc(cl::sycl::range<1>(shared_mem), cgh);
@@ -377,7 +377,7 @@ struct TensorEvaluator<const TensorConvolutionOp<Indices, InputArgType, KernelAr
           size_t range_x, GRange_x, tileSize_x, range_y, GRange_y, tileSize_y, range_z, GRange_z, tileSize_z;
           m_device.parallel_for_setup(numX, numY, numP, tileSize_x, tileSize_y, tileSize_z, range_x, range_y, range_z, GRange_x, GRange_y, GRange_z );
           const size_t shared_mem =(tileSize_x +kernel_size_x -1)*(tileSize_y +kernel_size_y -1) * tileSize_z;
-          assert(static_cast<unsigned long>(shared_mem) <= m_device.sharedMemPerBlock());
+          gpu_assert(static_cast<unsigned long>(shared_mem) <= m_device.sharedMemPerBlock());
           auto global_range=cl::sycl::range<3>(GRange_x, GRange_y, GRange_z);  // global range
           auto local_range=cl::sycl::range<3>(tileSize_x, tileSize_y, tileSize_z);  // local range
           InputLocalAcc local_acc(cl::sycl::range<1>(shared_mem), cgh);
@@ -408,7 +408,7 @@ struct TensorEvaluator<const TensorConvolutionOp<Indices, InputArgType, KernelAr
           size_t range_x, GRange_x, tileSize_x, range_y, GRange_y, tileSize_y, range_z, GRange_z, tileSize_z;
           m_device.parallel_for_setup(numX, numY, numZ, tileSize_x, tileSize_y, tileSize_z, range_x, range_y, range_z, GRange_x, GRange_y, GRange_z );
           const size_t shared_mem =(tileSize_x +kernel_size_x -1)*(tileSize_y +kernel_size_y -1) * (tileSize_z +kernel_size_y -1);
-          assert(static_cast<unsigned long>(shared_mem) <= m_device.sharedMemPerBlock());
+          gpu_assert(static_cast<unsigned long>(shared_mem) <= m_device.sharedMemPerBlock());
           auto global_range=cl::sycl::range<3>(GRange_x, GRange_y, GRange_z);  // global range
           auto local_range=cl::sycl::range<3>(tileSize_x, tileSize_y, tileSize_z);  // local range
           InputLocalAcc local_acc(cl::sycl::range<1>(shared_mem), cgh);
-- 
cgit v1.2.3