diff --git a/src/Grid/cuda/cuda_grid_gpu_funcs.cuh b/src/Grid/cuda/cuda_grid_gpu_funcs.cuh index ba1d1d4f2a3bc6e39c737ab271694cf7a96c6ff4..5560c3fe8df6693bb1ccdef7a8ea446d02c7706b 100644 --- a/src/Grid/cuda/cuda_grid_gpu_funcs.cuh +++ b/src/Grid/cuda/cuda_grid_gpu_funcs.cuh @@ -94,12 +94,12 @@ __global__ void copy_ndim_grid_device(grid_type src, grid_type dst) #endif -template<bool inte_or_lin,unsigned int dim, typename T> +template<bool inte_or_lin, typename base_grid, unsigned int dim, typename T> struct grid_toKernelImpl { - template<typename grid_type> static grid_gpu_ker<dim,T,memory_traits_lin,typename grid_type::linearizer_type> toKernel(grid_type & gc) + template<typename grid_type> static base_grid toKernel(grid_type & gc) { - grid_gpu_ker<dim,T,memory_traits_lin,typename grid_type::linearizer_type> g(gc.getGrid()); + /*grid_gpu_ker<dim,T,memory_traits_lin,typename grid_type::linearizer_type>*/base_grid g(gc.getGrid()); g.get_data_().disable_manage_memory(); g.get_data_().mem = gc.get_internal_data_().mem; @@ -112,12 +112,12 @@ struct grid_toKernelImpl } }; -template<unsigned int dim, typename T> -struct grid_toKernelImpl<true,dim,T> +template<typename base_grid, unsigned int dim, typename T> +struct grid_toKernelImpl<true,base_grid,dim,T> { - template<typename grid_type> static grid_gpu_ker<dim,T,memory_traits_inte, typename grid_type::linearizer_type> toKernel(grid_type & gc) + template<typename grid_type> static base_grid toKernel(grid_type & gc) { - grid_gpu_ker<dim,T,memory_traits_inte, typename grid_type::linearizer_type> g(gc.getGrid()); + /*grid_gpu_ker<dim,T,memory_traits_inte, typename grid_type::linearizer_type>*/ base_grid g(gc.getGrid()); copy_switch_memory_c_no_cpy<typename std::remove_reference<decltype(gc.get_internal_data_())>::type, typename std::remove_reference<decltype(g.get_data_())>::type> cp_mc(gc.get_internal_data_(),g.get_data_()); diff --git a/src/Grid/grid_base_implementation.hpp b/src/Grid/grid_base_implementation.hpp index 6143093d9b1e042ed035eb7d42ace5ec8847727d..9428e1be8755dc79876376573373105d9759742f 100644 --- a/src/Grid/grid_base_implementation.hpp +++ b/src/Grid/grid_base_implementation.hpp @@ -620,7 +620,7 @@ private: #if defined(CUDIFY_USE_SEQUENTIAL) || defined(CUDIFY_USE_OPENMP) - base_gpu = grid_toKernelImpl<is_layout_inte<layout_base<T_>>::value,dim,T_>::toKernel(*this); + base_gpu = grid_toKernelImpl<is_layout_inte<layout_base<T_>>::value,grid_gpu_ker<dim,T_,layout_base,linearizer_type>,dim,T_>::toKernel(*this); #endif } @@ -705,8 +705,8 @@ public: #if defined(CUDIFY_USE_SEQUENTIAL) || defined(CUDIFY_USE_OPENMP) - base_gpu = grid_toKernelImpl<is_layout_inte<layout_base<T_>>::value,dim,T_>::toKernel(*this); - g.base_gpu = grid_toKernelImpl<is_layout_inte<layout_base<T_>>::value,dim,T_>::toKernel(g); + base_gpu = grid_toKernelImpl<is_layout_inte<layout_base<T_>>::value,grid_gpu_ker<dim,T_,layout_base,linearizer_type>,dim,T_>::toKernel(*this); + g.base_gpu = grid_toKernelImpl<is_layout_inte<layout_base<T_>>::value,grid_gpu_ker<dim,T_,layout_base,linearizer_type>,dim,T_>::toKernel(g); #endif @@ -726,7 +726,7 @@ public: #if defined(CUDIFY_USE_SEQUENTIAL) || defined(CUDIFY_USE_OPENMP) - base_gpu = grid_toKernelImpl<is_layout_inte<layout_base<T_>>::value,dim,T_>::toKernel(*this); + base_gpu = grid_toKernelImpl<is_layout_inte<layout_base<T_>>::value,grid_gpu_ker<dim,T,layout_base,linearizer_type>,dim,T_>::toKernel(*this); #endif @@ -850,7 +850,7 @@ public: #if defined(CUDIFY_USE_SEQUENTIAL) || defined(CUDIFY_USE_OPENMP) - base_gpu = grid_toKernelImpl<is_layout_inte<layout_base<T_>>::value,dim,T_>::toKernel(*this); + base_gpu = grid_toKernelImpl<is_layout_inte<layout_base<T_>>::value,grid_gpu_ker<dim,T_,layout_base,linearizer_type>,dim,T_>::toKernel(*this); #endif @@ -893,7 +893,7 @@ public: #if defined(CUDIFY_USE_SEQUENTIAL) || defined(CUDIFY_USE_OPENMP) - base_gpu = grid_toKernelImpl<is_layout_inte<layout_base<T_>>::value,dim,T_>::toKernel(*this); + base_gpu = grid_toKernelImpl<is_layout_inte<layout_base<T_>>::value,grid_gpu_ker<dim,T_,layout_base,linearizer_type>,dim,T_>::toKernel(*this); #endif } @@ -922,7 +922,7 @@ public: #if defined(CUDIFY_USE_SEQUENTIAL) || defined(CUDIFY_USE_OPENMP) - base_gpu = grid_toKernelImpl<is_layout_inte<layout_base<T_>>::value,dim,T_>::toKernel(*this); + base_gpu = grid_toKernelImpl<is_layout_inte<layout_base<T_>>::value,grid_gpu_ker<dim,T_,layout_base,linearizer_type>,dim,T_>::toKernel(*this); #endif } @@ -1519,7 +1519,7 @@ public: #if defined(CUDIFY_USE_SEQUENTIAL) || defined(CUDIFY_USE_OPENMP) - base_gpu = grid_toKernelImpl<is_layout_inte<layout_base<T_>>::value,dim,T_>::toKernel(*this); + base_gpu = grid_toKernelImpl<is_layout_inte<layout_base<T_>>::value,grid_gpu_ker<dim,T_,layout_base,linearizer_type>,dim,T_>::toKernel(*this); #endif } @@ -1549,8 +1549,8 @@ public: #if defined(CUDIFY_USE_SEQUENTIAL) || defined(CUDIFY_USE_OPENMP) - base_gpu = grid_toKernelImpl<is_layout_inte<layout_base<T_>>::value,dim,T_>::toKernel(*this); - grid.base_gpu = grid_toKernelImpl<is_layout_inte<layout_base<T_>>::value,dim,T_>::toKernel(grid); + base_gpu = grid_toKernelImpl<is_layout_inte<layout_base<T_>>::value,grid_gpu_ker<dim,T_,layout_base,linearizer_type>,dim,T_>::toKernel(*this); + grid.base_gpu = grid_toKernelImpl<is_layout_inte<layout_base<T_>>::value,grid_gpu_ker<dim,T_,layout_base,linearizer_type>,dim,T_>::toKernel(grid); #endif } diff --git a/src/SparseGridGpu/SparseGridGpu_kernels.cuh b/src/SparseGridGpu/SparseGridGpu_kernels.cuh index 7b7d5c10d6b36b4aa836d069d8850dac56821210..683dc776c821b29c9bc6815fd5433d5ce236cf7a 100644 --- a/src/SparseGridGpu/SparseGridGpu_kernels.cuh +++ b/src/SparseGridGpu/SparseGridGpu_kernels.cuh @@ -966,7 +966,7 @@ namespace SparseGridGpuKernels auto dataBlockLoad = dataBuffer.get(dataBlockPos); // Avoid binary searches as much as possible // todo: Add management of RED-BLACK stencil application! :) - const unsigned int dataBlockId = indexBuffer.template get<pIndex>(dataBlockPos); + const auto dataBlockId = indexBuffer.template get<pIndex>(dataBlockPos); grid_key_dx<dim, int> pointCoord = sparseGrid.getCoord(dataBlockId * blockSize + offset); unsigned char curMask;