Skip to content

Commit 1ab137b

Browse files
ericcanoericcano
committed
Added SoA stores and views and a test
The test validates multiple multiple SoA use cases, host and device side. It validates buffer sharing for multiple SoA. An an Alpaka native sub-buffer operation (device side memset) was used. This type of operations can be used in the future to copy partial columns.
1 parent 2cf0ae8 commit 1ab137b

4 files changed

Lines changed: 1385 additions & 0 deletions

File tree

‎src/alpaka/DataFormats/SoACommon.h‎

Lines changed: 183 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,183 @@
1+
/*
2+
* Definitions of SoA common parameters for SoA class generators
3+
*/
4+
5+
#ifndef DataStructures_SoACommon_h
6+
#define DataStructures_SoACommon_h
7+
8+
#include "boost/preprocessor.hpp"
9+
#include <Eigen/Core>
10+
11+
// CUDA attributes
12+
#ifdef __CUDACC__
13+
#define SOA_HOST_ONLY __host__
14+
#define SOA_DEVICE_ONLY __device__
15+
#define SOA_HOST_DEVICE __host__ __device__
16+
#define SOA_HOST_DEVICE_INLINE __host__ __device__ __forceinline__
17+
#define SOA_DEVICE_RESTRICT __restrict__
18+
#else
19+
#define SOA_HOST_ONLY
20+
#define SOA_DEVICE_ONLY
21+
#define SOA_HOST_DEVICE
22+
#define SOA_HOST_DEVICE_INLINE inline
23+
#define SOA_DEVICE_RESTRICT
24+
#endif
25+
26+
#if defined(__CUDACC__) && defined(__CUDA_ARCH__)
27+
// Read a pointer content via read-only (non coherent) cache.
28+
#define LOAD_INCOHERENT(A) __ldg(A)
29+
#define LOAD_STREAMED(A) __ldcs(A)
30+
#define STORE_STREAMED(A, V) __stcs(A, V)
31+
#else
32+
#define LOAD_INCOHERENT(A) *(A)
33+
#define LOAD_STREAMED(A) *(A)
34+
#define STORE_STREAMED(A, V) *(A) = (V)
35+
#endif
36+
37+
// compile-time sized SoA
38+
39+
// Helper template managing the value within it column
40+
template<typename T>
41+
class SoAValue {
42+
public:
43+
SOA_HOST_DEVICE_INLINE SoAValue(size_t i, T * col): idx_(i), col_(col) {}
44+
/* SOA_HOST_DEVICE_INLINE operator T&() { return col_[idx_]; } */
45+
SOA_HOST_DEVICE_INLINE T& operator() () { return col_[idx_]; }
46+
SOA_HOST_DEVICE_INLINE T operator() () const { return *(col_ + idx_); }
47+
SOA_HOST_DEVICE_INLINE T* operator& () { return &col_[idx_]; }
48+
SOA_HOST_DEVICE_INLINE const T* operator& () const { return &col_[idx_]; }
49+
template <typename T2>
50+
SOA_HOST_DEVICE_INLINE T& operator= (const T2& v) { return col_[idx_] = v; }
51+
typedef T valueType;
52+
static constexpr auto valueSize = sizeof(T);
53+
private:
54+
size_t idx_;
55+
T *col_;
56+
};
57+
58+
// Helper template managing the value within it column
59+
template<typename T>
60+
class SoAConstValue {
61+
public:
62+
SOA_HOST_DEVICE_INLINE SoAConstValue(size_t i, const T * col): idx_(i), col_(col) {}
63+
/* SOA_HOST_DEVICE_INLINE operator T&() { return col_[idx_]; } */
64+
SOA_HOST_DEVICE_INLINE T operator() () const { return *(col_ + idx_); }
65+
SOA_HOST_DEVICE_INLINE const T* operator& () const { return &col_[idx_]; }
66+
typedef T valueType;
67+
static constexpr auto valueSize = sizeof(T);
68+
private:
69+
size_t idx_;
70+
const T *col_;
71+
};
72+
73+
74+
// Helper template managing the value within it column
75+
template<class C>
76+
class SoAEigenValue {
77+
public:
78+
typedef C Type;
79+
typedef Eigen::Map<C, 0, Eigen::InnerStride<Eigen::Dynamic>> MapType;
80+
typedef Eigen::Map<const C, 0, Eigen::InnerStride<Eigen::Dynamic>> CMapType;
81+
SOA_HOST_DEVICE_INLINE SoAEigenValue(size_t i, typename C::Scalar * col, size_t stride):
82+
val_(col + i, C::RowsAtCompileTime, C::ColsAtCompileTime,
83+
Eigen::InnerStride<Eigen::Dynamic>(stride)),
84+
crCol_(col),
85+
cVal_(crCol_ + i, C::RowsAtCompileTime, C::ColsAtCompileTime,
86+
Eigen::InnerStride<Eigen::Dynamic>(stride)),
87+
stride_(stride) {}
88+
SOA_HOST_DEVICE_INLINE MapType& operator() () { return val_; }
89+
SOA_HOST_DEVICE_INLINE const CMapType& operator() () const { return cVal_; }
90+
SOA_HOST_DEVICE_INLINE operator C() { return val_; }
91+
SOA_HOST_DEVICE_INLINE operator const C() const { return cVal_; }
92+
SOA_HOST_DEVICE_INLINE C* operator& () { return &val_; }
93+
SOA_HOST_DEVICE_INLINE const C* operator& () const { return &cVal_; }
94+
template <class C2>
95+
SOA_HOST_DEVICE_INLINE MapType& operator= (const C2& v) { return val_ = v; }
96+
typedef typename C::Scalar ValueType;
97+
static constexpr auto valueSize = sizeof(C::Scalar);
98+
SOA_HOST_DEVICE_INLINE size_t stride() { return stride_; }
99+
template<typename OtherDerived>
100+
typename Eigen::MatrixBase<C>::template cross_product_return_type<OtherDerived>::type
101+
SOA_HOST_DEVICE_INLINE cross(const Eigen::MatrixBase<OtherDerived>& other) const { return cVal_.cross(other); }
102+
103+
template<typename OtherType>
104+
typename Eigen::MatrixBase<C>::template cross_product_return_type<typename OtherType::MapType>::type
105+
SOA_HOST_DEVICE_INLINE cross(const OtherType& other) const { return cVal_.cross(other.cVal_); }
106+
107+
private:
108+
MapType val_;
109+
const typename C::Scalar * __restrict__ crCol_;
110+
CMapType cVal_;
111+
size_t stride_;
112+
};
113+
114+
// Helper template to avoid commas in macro
115+
template<class C>
116+
struct EigenConstMapMaker {
117+
typedef Eigen::Map<const C, 0, Eigen::InnerStride<Eigen::Dynamic>> Type;
118+
class DataHolder {
119+
public:
120+
DataHolder(const typename C::Scalar * data): data_(data) {}
121+
EigenConstMapMaker::Type withStride(size_t stride) {
122+
return EigenConstMapMaker::Type(data_, C::RowsAtCompileTime, C::ColsAtCompileTime,
123+
Eigen::InnerStride<Eigen::Dynamic>(stride));
124+
}
125+
private:
126+
const typename C::Scalar * const data_;
127+
};
128+
static DataHolder withData(const typename C::Scalar * data) {
129+
return DataHolder(data);
130+
}
131+
};
132+
133+
// Helper function to compute aligned size
134+
inline size_t alignSize(size_t size, size_t alignment = 128) {
135+
if (size)
136+
return ((size - 1) / alignment + 1) * alignment;
137+
else
138+
return 0;
139+
}
140+
141+
/* declare "scalars" (one value shared across the whole SoA) and "columns" (one value per element) */
142+
#define _VALUE_TYPE_SCALAR 0
143+
#define _VALUE_TYPE_COLUMN 1
144+
#define _VALUE_TYPE_EIGEN_COLUMN 2
145+
146+
enum class SoAColumnType {
147+
scalar = _VALUE_TYPE_SCALAR,
148+
column = _VALUE_TYPE_COLUMN,
149+
eigen = _VALUE_TYPE_EIGEN_COLUMN
150+
};
151+
152+
#define SoA_scalar(TYPE, NAME) (_VALUE_TYPE_SCALAR, TYPE, NAME)
153+
#define SoA_column(TYPE, NAME) (_VALUE_TYPE_COLUMN, TYPE, NAME)
154+
#define SoA_eigenColumn(TYPE, NAME) (_VALUE_TYPE_EIGEN_COLUMN, TYPE, NAME)
155+
156+
/* Iterate on the macro MACRO and return the result as a comma separated list */
157+
#define _ITERATE_ON_ALL_COMMA(MACRO, DATA, ...) \
158+
BOOST_PP_TUPLE_ENUM( \
159+
BOOST_PP_SEQ_TO_TUPLE( \
160+
_ITERATE_ON_ALL(MACRO, DATA, __VA_ARGS__) \
161+
) \
162+
)
163+
/* Iterate MACRO on all elements */
164+
#define _ITERATE_ON_ALL(MACRO, DATA, ...) \
165+
BOOST_PP_SEQ_FOR_EACH(MACRO, DATA, \
166+
BOOST_PP_VARIADIC_TO_SEQ(__VA_ARGS__) \
167+
)
168+
169+
/* Switch on macros depending on scalar / column type */
170+
#define _SWITCH_ON_TYPE(VALUE_TYPE, IF_SCALAR, IF_COLUMN, IF_EIGEN_COLUMN) \
171+
BOOST_PP_IF(BOOST_PP_EQUAL(VALUE_TYPE, _VALUE_TYPE_SCALAR), \
172+
IF_SCALAR, \
173+
BOOST_PP_IF(BOOST_PP_EQUAL(VALUE_TYPE, _VALUE_TYPE_COLUMN), \
174+
IF_COLUMN, \
175+
BOOST_PP_IF(BOOST_PP_EQUAL(VALUE_TYPE, _VALUE_TYPE_EIGEN_COLUMN), \
176+
IF_EIGEN_COLUMN, \
177+
BOOST_PP_EMPTY() \
178+
) \
179+
) \
180+
)
181+
182+
#endif // ndef DataStructures_SoACommon_h
183+

0 commit comments

Comments
 (0)