Homec4science
Diffusion tamaas (master) Browse

tamaas/src/gpu/loops/loop_cuda.hh
b3c9bf19b545master

Last Change
Raw File

loop_cuda.hh
View Options

/**
* @file
*
* @author Lucas Frérot <lucas.frerot@epfl.ch>
*
* @section LICENSE
*
* Copyright (©) 2017 EPFL (Ecole Polytechnique Fédérale de
* Lausanne) Laboratory (LSMS - Laboratoire de Simulation en Mécanique des
* Solides)
*
* Tamaas is free software: you can redistribute it and/or modify it under the
* terms of the GNU Lesser General Public License as published by the Free
* Software Foundation, either version 3 of the License, or (at your option) any
* later version.
*
* Tamaas is distributed in the hope that it will be useful, but WITHOUT ANY
* WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS FOR
* A PARTICULAR PURPOSE. See the GNU Lesser General Public License for more
* details.
*
* You should have received a copy of the GNU Lesser General Public License
* along with Tamaas. If not, see <http://www.gnu.org/licenses/>.
*
*/
/* -------------------------------------------------------------------------- */
#ifndef __LOOP_CUDA_HH__
#define __LOOP_CUDA_HH__
/* -------------------------------------------------------------------------- */
#include "tamaas.hh"
#include "loops/loop_utils.hh"
#include <cuda_runtime_api.h>
__BEGIN_TAMAAS__
namespace detail {
#define GRID_STRIDED_LOOP(operation, size) do { \
const UInt index = blockIdx.x * blockDim.x + threadIdx.x; \
const UInt stride = blockDim.x * gridDim.x; \
for (UInt i = index ; i < (size) ; i += stride) { \
(operation); \
}} while(0) \
template <bool strided, typename Functor, typename... T>
__global__ void exec_loop_kernel(Functor& func, UInt size, T... data_pointers) {
//GRID_STRIDED_LOOP(func(data_pointers[i]...), size);
UInt index = blockIdx.x * blockDim.x + threadIdx.x;
UInt stride = blockDim.x * gridDim.x;
for (UInt i = index ; i < size ; i += stride) {
func(data_pointers[i]...);
}
}
template <bool strided, typename Functor, typename... Grids>
void exec_loop_cuda(Functor&& func, Grids&&... grids) {
Int device = 0, threads_per_block = 0;
cudaGetDevice(&device);
cudaDeviceGetAttribute(&threads_per_block, cudaDevAttrMaxThreadsPerBlock, device);
const UInt loop_size = loopSize<strided>(grids...);
UInt blocks = loop_size / threads_per_block;
blocks += (threads_per_block * blocks == loop_size)? 0 : 1; // make sure we cover all data
//exec_loop_kernel<strided><<<blocks, threads_per_block>>>(func, loop_size, grids.getInternalData()...);
exec_loop_kernel<strided><<<1, 1>>>(func, loop_size, grids.getInternalData()...);
cudaDeviceSynchronize();
}
template <operation op, bool strided, typename Functor, typename... Grids>
auto reduce_loop_cuda(Functor&& func, Grids&&... grids) -> decltype(func(grids(0)...)) {
return func(grids(0)...);
}
#undef GRID_STRIDED_LOOP
} // namespace detail
__END_TAMAAS__
#endif // __LOOP_CUDA_HH__
c4science · Help