Page Menu
Home
c4science
Search
Configure Global Search
Log In
Files
F85190957
loop.hh
No One
Temporary
Actions
Download File
Edit File
Delete File
View Transforms
Subscribe
Mute Notifications
Award Token
Subscribers
None
File Metadata
Details
File Info
Storage
Attached
Created
Fri, Sep 27, 10:07
Size
8 KB
Mime Type
text/x-c++
Expires
Sun, Sep 29, 10:07 (2 d)
Engine
blob
Format
Raw Data
Handle
21119843
Attached To
rTAMAAS tamaas
loop.hh
View Options
/**
* @file
* @section LICENSE
*
* Copyright (©) 2016-19 EPFL (École Polytechnique Fédérale de Lausanne),
* Laboratory (LSMS - Laboratoire de Simulation en Mécanique des Solides)
*
* This program is free software: you can redistribute it and/or modify
* it under the terms of the GNU Affero General Public License as published
* by the Free Software Foundation, either version 3 of the License, or
* (at your option) any later version.
*
* This program 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 Affero General Public License for more details.
*
* You should have received a copy of the GNU Affero General Public License
* along with this program. If not, see <https://www.gnu.org/licenses/>.
*
*/
/* -------------------------------------------------------------------------- */
#ifndef LOOP_HH
#define LOOP_HH
/* -------------------------------------------------------------------------- */
#include "loops/apply.hh"
#include "loops/loop_utils.hh"
#include "ranges.hh"
#include "tamaas.hh"
#include <thrust/execution_policy.h>
#include <thrust/for_each.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/transform_reduce.h>
#include <thrust/tuple.h>
#include <thrust/version.h>
#include <type_traits>
#include <utility>
namespace
tamaas
{
template
<
typename
T
>
struct
is_policy
:
std
::
false_type
{};
template
<>
struct
is_policy
<
thrust
::
detail
::
host_t
>
:
std
::
true_type
{};
template
<>
struct
is_policy
<
thrust
::
detail
::
device_t
>
:
std
::
true_type
{};
template
<>
struct
is_policy
<
const
thrust
::
detail
::
host_t
>
:
std
::
true_type
{};
template
<>
struct
is_policy
<
const
thrust
::
detail
::
device_t
>
:
std
::
true_type
{};
template
<>
struct
is_policy
<
const
thrust
::
detail
::
host_t
&>
:
std
::
true_type
{};
template
<>
struct
is_policy
<
const
thrust
::
detail
::
device_t
&>
:
std
::
true_type
{};
/**
* @brief Singleton class for automated loops using lambdas
* This class is sweet candy :) It provides abstraction of the paralelism
* paradigm used in loops and allows simple and less error-prone loop syntax,
* with minimum boiler plate. I love it <3
*/
class
Loop
{
public
:
/// Backends enumeration
enum
backend
{
omp
,
///< [OpenMP](http://www.openmp.org/specifications/) backend
cuda
,
///< [Cuda](http://docs.nvidia.com/cuda/index.html) backend
};
/// Helper class to count iterations within lambda-loop
template
<
typename
T
>
class
arange
{
public
:
using
it_type
=
thrust
::
counting_iterator
<
T
>
;
using
reference
=
typename
it_type
::
reference
;
arange
(
T
start
,
T
size
)
:
start
(
start
),
range_size
(
size
)
{}
it_type
begin
(
UInt
/*i*/
=
1
)
const
{
return
it_type
(
T
(
start
));
}
it_type
end
(
UInt
/*i*/
=
1
)
const
{
return
it_type
(
range_size
);
}
UInt
getNbComponents
()
const
{
return
1
;
}
private
:
T
start
,
range_size
;
};
template
<
typename
T
>
static
arange
<
T
>
range
(
T
size
)
{
return
arange
<
T
>
(
0
,
size
);
}
template
<
typename
T
,
typename
U
>
static
arange
<
T
>
range
(
U
start
,
T
size
)
{
return
arange
<
T
>
(
start
,
size
);
}
private
:
/// Replacement for thrust::transform_iterator which copies values
template
<
typename
Iterator
,
typename
Functor
,
typename
Value
>
class
transform_iterator
:
public
thrust
::
iterator_adaptor
<
transform_iterator
<
Iterator
,
Functor
,
Value
>
,
Iterator
,
Value
,
thrust
::
use_default
,
thrust
::
use_default
,
Value
>
{
Functor
func
;
public
:
using
super_t
=
thrust
::
iterator_adaptor
<
transform_iterator
<
Iterator
,
Functor
,
Value
>
,
Iterator
,
Value
,
thrust
::
use_default
,
thrust
::
use_default
,
Value
>
;
__host__
__device__
transform_iterator
(
const
Iterator
&
x
,
const
Functor
&
func
)
:
super_t
(
x
),
func
(
func
)
{}
friend
class
thrust
::
iterator_core_access
;
private
:
__host__
__device__
auto
dereference
()
const
->
decltype
(
func
(
*
this
->
base
()))
{
return
func
(
*
this
->
base
());
}
};
public
:
/// Loop functor over ranges
template
<
typename
Functor
,
typename
...
Ranges
>
static
auto
loop
(
Functor
&&
func
,
Ranges
&&
...
ranges
)
->
typename
std
::
enable_if
<
not
is_policy
<
Functor
>::
value
,
void
>::
type
{
loopImpl
(
thrust
::
device
,
std
::
forward
<
Functor
>
(
func
),
std
::
forward
<
Ranges
>
(
ranges
)...);
}
/// Loop over ranges with non-default policy
template
<
typename
DerivedPolicy
,
typename
Functor
,
typename
...
Ranges
>
static
void
loop
(
const
thrust
::
execution_policy
<
DerivedPolicy
>&
policy
,
Functor
&&
func
,
Ranges
&&
...
ranges
)
{
loopImpl
(
policy
,
std
::
forward
<
Functor
>
(
func
),
std
::
forward
<
Ranges
>
(
ranges
)...);
}
private
:
template
<
typename
T
>
using
reference_type
=
typename
std
::
decay
<
T
>::
type
::
reference
;
public
:
/// Reduce functor over ranges
template
<
operation
op
,
typename
Functor
,
typename
...
Ranges
>
static
auto
reduce
(
Functor
&&
func
,
Ranges
&&
...
ranges
)
->
typename
std
::
enable_if
<
not
is_policy
<
Functor
>::
value
,
decltype
(
func
(
std
::
declval
<
reference_type
<
Ranges
>>
()...))
>::
type
{
return
reduceImpl
<
op
>
(
thrust
::
device
,
std
::
forward
<
Functor
>
(
func
),
std
::
forward
<
Ranges
>
(
ranges
)...);
}
/// Reduce over ranges with non-default policy
template
<
operation
op
,
typename
DerivedPolicy
,
typename
Functor
,
typename
...
Ranges
>
static
auto
reduce
(
const
thrust
::
execution_policy
<
DerivedPolicy
>&
policy
,
Functor
&&
func
,
Ranges
&&
...
ranges
)
->
decltype
(
func
(
std
::
declval
<
reference_type
<
Ranges
>>
()...))
{
return
reduceImpl
<
op
>
(
policy
,
std
::
forward
<
Functor
>
(
func
),
std
::
forward
<
Ranges
>
(
ranges
)...);
}
private
:
/// Loop over ranges and apply functor
template
<
typename
DerivedPolicy
,
typename
Functor
,
typename
...
Ranges
>
static
void
loopImpl
(
const
thrust
::
execution_policy
<
DerivedPolicy
>&
policy
,
Functor
&&
func
,
Ranges
&&
...
ranges
);
/// Loop over ranges, apply functor and reduce result
template
<
operation
op
,
typename
DerivedPolicy
,
typename
Functor
,
typename
...
Ranges
>
static
auto
reduceImpl
(
const
thrust
::
execution_policy
<
DerivedPolicy
>&
policy
,
Functor
&&
func
,
Ranges
&&
...
ranges
)
->
decltype
(
func
(
std
::
declval
<
reference_type
<
Ranges
>>
()...));
/// Constructor
Loop
()
=
delete
;
};
/* -------------------------------------------------------------------------- */
/* Template implementation */
/* -------------------------------------------------------------------------- */
template
<
typename
DerivedPolicy
,
typename
Functor
,
typename
...
Ranges
>
void
Loop
::
loopImpl
(
const
thrust
::
execution_policy
<
DerivedPolicy
>&
policy
,
Functor
&&
func
,
Ranges
&&
...
ranges
)
{
auto
begin
=
thrust
::
make_zip_iterator
(
thrust
::
make_tuple
(
ranges
.
begin
()...));
auto
end
=
thrust
::
make_zip_iterator
(
thrust
::
make_tuple
(
ranges
.
end
()...));
checkLoopSize
(
ranges
...);
thrust
::
for_each
(
policy
,
begin
,
end
,
detail
::
ApplyFunctor
<
Functor
>
(
std
::
forward
<
Functor
>
(
func
)));
#if (defined(USE_CUDA) && THRUST_VERSION < 100904)
cudaDeviceSynchronize
();
#endif
}
/* -------------------------------------------------------------------------- */
template
<
operation
op
,
typename
DerivedPolicy
,
typename
Functor
,
typename
...
Ranges
>
auto
Loop
::
reduceImpl
(
const
thrust
::
execution_policy
<
DerivedPolicy
>&
policy
,
Functor
&&
func
,
Ranges
&&
...
ranges
)
->
decltype
(
func
(
std
::
declval
<
reference_type
<
Ranges
>>
()...))
{
using
return_type
=
decltype
(
func
(
std
::
declval
<
reference_type
<
Ranges
>>
()...));
using
apply_type
=
detail
::
ApplyFunctor
<
Functor
,
return_type
>
;
checkLoopSize
(
ranges
...);
auto
applier
=
apply_type
(
func
);
detail
::
reduction_helper
<
op
,
return_type
>
red_helper
;
auto
begin_zip
=
thrust
::
make_zip_iterator
(
thrust
::
make_tuple
(
ranges
.
begin
()...));
auto
end_zip
=
thrust
::
make_zip_iterator
(
thrust
::
make_tuple
(
ranges
.
end
()...));
transform_iterator
<
decltype
(
begin_zip
),
apply_type
,
return_type
>
begin
(
begin_zip
,
applier
);
transform_iterator
<
decltype
(
end_zip
),
apply_type
,
return_type
>
end
(
end_zip
,
applier
);
auto
result
=
thrust
::
reduce
(
policy
,
begin
,
end
,
red_helper
.
init
(),
red_helper
);
#if (defined(USE_CUDA) && THRUST_VERSION < 100904)
cudaDeviceSynchronize
();
#endif
return
result
;
}
}
// namespace tamaas
#endif
// LOOP_HH
Event Timeline
Log In to Comment