mirror of
https://gitlab.com/libeigen/eigen.git
synced 2025-08-13 20:26:03 +08:00
Removing unused variables
This commit is contained in:
parent
837fdbdcb2
commit
77cc4d06c7
@ -326,7 +326,7 @@ struct TensorEvaluator<const TensorConvolutionOp<Indices, InputArgType, KernelAr
|
||||
// extract input functor list
|
||||
InputFunctorExpr input_functors = Eigen::TensorSycl::internal::extractFunctors(m_inputImpl);
|
||||
|
||||
const unsigned long maxSharedMem = m_device.sharedMemPerBlock(); // sycl localmemory size
|
||||
|
||||
m_device.sycl_queue().submit([&](cl::sycl::handler &cgh) {
|
||||
|
||||
typedef cl::sycl::accessor<CoeffReturnType, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local> InputLocalAcc;
|
||||
@ -348,7 +348,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) <= maxSharedMem);
|
||||
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);
|
||||
@ -373,7 +373,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) <= maxSharedMem);
|
||||
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);
|
||||
@ -404,7 +404,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) <= maxSharedMem);
|
||||
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);
|
||||
|
@ -386,7 +386,7 @@ struct SyclDevice {
|
||||
return 2;//sycl_queue().get_device(). template get_info<cl::sycl::info::device::max_work_group_size>();
|
||||
// return stream_->deviceProperties().maxThreadsPerMultiProcessor;
|
||||
}
|
||||
EIGEN_STRONG_INLINE int sharedMemPerBlock() const {
|
||||
EIGEN_STRONG_INLINE size_t sharedMemPerBlock() const {
|
||||
return sycl_queue().get_device(). template get_info<cl::sycl::info::device::local_mem_size>();
|
||||
// return stream_->deviceProperties().sharedMemPerBlock;
|
||||
}
|
||||
|
Loading…
x
Reference in New Issue
Block a user