Stokhos Package Browser (Single Doxygen Collection) Version of the Day
Loading...
Searching...
No Matches
Stokhos_Cuda_WarpShuffle.hpp
Go to the documentation of this file.
1// @HEADER
2// ***********************************************************************
3//
4// Stokhos Package
5// Copyright (2009) Sandia Corporation
6//
7// Under terms of Contract DE-AC04-94AL85000, there is a non-exclusive
8// license for use of this work by or on behalf of the U.S. Government.
9//
10// Redistribution and use in source and binary forms, with or without
11// modification, are permitted provided that the following conditions are
12// met:
13//
14// 1. Redistributions of source code must retain the above copyright
15// notice, this list of conditions and the following disclaimer.
16//
17// 2. Redistributions in binary form must reproduce the above copyright
18// notice, this list of conditions and the following disclaimer in the
19// documentation and/or other materials provided with the distribution.
20//
21// 3. Neither the name of the Corporation nor the names of the
22// contributors may be used to endorse or promote products derived from
23// this software without specific prior written permission.
24//
25// THIS SOFTWARE IS PROVIDED BY SANDIA CORPORATION "AS IS" AND ANY
26// EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
27// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
28// PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL SANDIA CORPORATION OR THE
29// CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
30// EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
31// PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
32// PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
33// LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
34// NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
35// SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
36//
37// Questions? Contact Eric T. Phipps (etphipp@sandia.gov).
38//
39// ***********************************************************************
40// @HEADER
41
42#ifndef STOKHOS_CUDA_WARP_SHUFFLE_HPP
43#define STOKHOS_CUDA_WARP_SHUFFLE_HPP
44
45#include "Kokkos_Core.hpp"
46
47#ifdef __CUDA_ARCH__
48# if (__CUDA_ARCH__ >= 300)
49# define HAVE_CUDA_SHUFFLE 1
50# else
51# define HAVE_CUDA_SHUFFLE 0
52# endif
53#else
54# define HAVE_CUDA_SHUFFLE 0
55#endif
56
57namespace Stokhos {
58
59template<typename Scalar>
60KOKKOS_INLINE_FUNCTION
61Scalar shfl_down(const Scalar &val, const int& delta, const int& width){
62 return Kokkos::shfl_down(val, delta, width);
63}
64
65template<typename Scalar>
66KOKKOS_INLINE_FUNCTION
67Scalar shfl_up(const Scalar &val, const int& delta, const int& width){
68 return Kokkos::shfl_up(val, delta, width);
69}
70
71template<typename Scalar>
72KOKKOS_INLINE_FUNCTION
73Scalar shfl_down(const Scalar &val, const int& delta, const int& width,
74 const int& mask){
75#ifdef __CUDA_ARCH__
76 return __shfl_down_sync(mask, val, delta, width);
77#else
78 (void) delta;
79 (void) width;
80 (void) mask;
81 return val;
82#endif
83}
84
85template<typename Scalar>
86KOKKOS_INLINE_FUNCTION
87Scalar shfl_up(const Scalar &val, const int& delta, const int& width,
88 const int& mask){
89#ifdef __CUDA_ARCH__
90 return __shfl_up_sync(mask, val, delta, width);
91#else
92 (void) delta;
93 (void) width;
94 (void) mask;
95 return val;
96#endif
97}
98
99KOKKOS_INLINE_FUNCTION
100void sync_warp(const int& mask) {
101#ifdef __CUDA_ARCH__
102 __syncwarp(mask);
103#endif
104}
105
106} // namespace Stokhos
107
108#endif /* #ifndef STOKHOS_CUDA_WARP_SHUFFLE_HPP */
expr val()
Top-level namespace for Stokhos classes and functions.
KOKKOS_INLINE_FUNCTION void sync_warp(const int &mask)
KOKKOS_INLINE_FUNCTION Scalar shfl_up(const Scalar &val, const int &delta, const int &width)
KOKKOS_INLINE_FUNCTION Scalar shfl_down(const Scalar &val, const int &delta, const int &width)