From 72f11b108ca85bd85f41857a4f4adf7452980d04 Mon Sep 17 00:00:00 2001
From: Behzad Safaei <iwia103h@alex1.nhr.fau.de>
Date: Wed, 5 Feb 2025 15:44:16 +0100
Subject: [PATCH] Enable device support in PairsAccessor (initial working
version)
---
runtime/math/MathTrait.h | 633 +++++++++++++++++++++++++++++++++
runtime/math/Vector3.hpp | 43 +++
runtime/pairs.hpp | 4 +
runtime/pairs_common.hpp | 10 +
src/pairs/code_gen/accessor.py | 369 +++++++++++++++++++
src/pairs/code_gen/cgen.py | 156 ++------
src/pairs/ir/types.py | 8 +-
7 files changed, 1097 insertions(+), 126 deletions(-)
create mode 100644 runtime/math/MathTrait.h
create mode 100644 runtime/math/Vector3.hpp
create mode 100644 src/pairs/code_gen/accessor.py
diff --git a/runtime/math/MathTrait.h b/runtime/math/MathTrait.h
new file mode 100644
index 0000000..44362c9
--- /dev/null
+++ b/runtime/math/MathTrait.h
@@ -0,0 +1,633 @@
+#pragma once
+
+//*************************************************************************************************
+// Includes
+//*************************************************************************************************
+
+#include <cstddef>
+
+namespace pairs {
+
+
+//=================================================================================================
+//
+// MATHEMATICAL TRAIT
+//
+//=================================================================================================
+
+//*************************************************************************************************
+/*!\class MathTrait
+ * \brief Base template for the MathTrait class.
+ * \ingroup math
+ *
+ * \section mathtrait_general General
+ *
+ * The MathTrait class template offers the possibility to select the resulting data type
+ * of a generic mathematical operation. In case of operations between built-in data types,
+ * the MathTrait class defines the more significant data type as the resulting data type.
+ * For this selection, signed data types are given a higher significance. In case of
+ * operations involving user-defined data types, the MathTrait template specifies the
+ * resulting data type of this operation.\n
+ * Specifying the resulting data type for a specific operation is done by specializing
+ * the MathTrait template for this particular type combination. In case a certain type
+ * combination is not defined in a MathTrait specialization, the base template is selected,
+ * which defines no resulting types and therefore stops the compilation process. Each
+ * specialization defines the data types \a HighType that represents the high-order data
+ * type of the two given data types and \a LowType that represents the low-order data type.
+ * Additionally, each specialization defines the types \a AddType, \a SubType, \a MultType
+ * and \a DivType, that represent the type of the resulting data type of the corresponding
+ * mathematical operation. The following example shows the specialization for operations
+ * between the double and the integer type:
+
+ \code
+ template<>
+ struct MathTrait< double, int >
+ {
+ typedef double HighType;
+ typedef int LowType;
+ typedef double AddType;
+ typedef double SubType;
+ typedef double MultType;
+ typedef double DivType;
+ };
+ \endcode
+
+ * Per default, the MathTrait template provides specializations for the following built-in
+ * data types:
+ *
+ * <ul>
+ * <li>integers</li>
+ * <ul>
+ * <li>unsigned char, signed char, char, wchar_t</li>
+ * <li>unsigned short, short</li>
+ * <li>unsigned int, int</li>
+ * <li>unsigned long, long</li>
+ * <li>std::size_t, std::ptrdiff_t (for certain 64-bit compilers)</li>
+ * </ul>
+ * <li>floating points</li>
+ * <ul>
+ * <li>float</li>
+ * <li>double</li>
+ * <li>long double</li>
+ * </ul>
+ * </ul>
+ *
+ *
+ * \n \section specializations Creating custom specializations
+ *
+ * It is possible to specialize the MathTrait template for additional user-defined data types.
+ * However, it is possible that a specific mathematical operation is invalid for the particular
+ * type combination. In this case, the INVALID_NUMERICAL_TYPE can be used to fill the missing
+ * type definition. The INVALID_NUMERICAL_TYPE represents the resulting data type of an invalid
+ * numerical operation. It is left undefined to stop the compilation process in case it is
+ * instantiated. The following example shows the specialization of the MathTrait template for
+ * Matrix3 and Vector3. In this case, only the multiplication between the matrix and the vector
+ * is a valid numerical operation. Therefore for all other types the INVALID_NUMERICAL_TYPE is
+ * used.
+
+ \code
+ template< typename T1, typename T2 >
+ struct MathTrait< Matrix3<T1>, Vector3<T2> >
+ {
+ typedef INVALID_NUMERICAL_TYPE HighType; // Invalid, no common high data type
+ typedef INVALID_NUMERICAL_TYPE LowType; // Invalid, no common low data type
+ typedef INVALID_NUMERICAL_TYPE AddType; // Invalid, cannot add a matrix and a vector
+ typedef INVALID_NUMERICAL_TYPE SubType; // Invalid, cannot subtract a vector from a matrix
+ typedef Vector3< typename MathTrait<T1,T2>::MultType > MultType; // Multiplication between a matrix and a vector
+ typedef INVALID_NUMERICAL_TYPE DivType; // Invalid, cannot divide a matrix by a vector
+ };
+ \endcode
+
+ * \n \section mathtrait_examples Examples
+ *
+ * The following example demonstrates the use of the MathTrait template, where depending on
+ * the two given data types the resulting data type is selected:
+
+ \code
+ template< typename T1, typename T2 > // The two generic types
+ typename MathTrait<T1,T2>::HighType // The resulting generic return type
+ add( T1 t1, T2 t2 ) //
+ { // The function 'add' returns the sum
+ return t1 + t2; // of the two given values
+ } //
+ \endcode
+
+ * Additionally, the specializations of the MathTrait template enable arithmetic operations
+ * between any combination of the supported data types:
+
+ \code
+ typedef Vector3< Matrix3< float > > VectorOfMatrices; // Vector of single-precision matrices
+ typedef Vector3< Vector3 < double > > VectorOfVectors; // Vector of double-precision vectors
+ typedef Vector3< double > VectorOfScalars; // Vector of double-precision scalars
+
+ VectorOfMatrices vm; // Setup of a vector of matrices
+ VectorOfVectors vv; // Setup of a vector of vectors
+
+ // Calculation of the scalar product between the two vectors. The resulting data type
+ // is a plain 3-dimensional vector of scalar values of type double.
+ VectorOfScalars res = vm * vv;
+ \endcode
+ */
+//*************************************************************************************************
+
+//strange but needed for compatibility reasons with visual studio compiler
+//backward compatibility to old PAIRS code
+template< typename T1, typename T2 >
+struct MathTrait
+{
+ using HighType = T1;
+ using LowType = T2;
+ using High = T1;
+ using Low = T2;
+};
+
+template< typename T>
+struct MathTrait< T, T >
+{
+ using HighType = T;
+ using LowType = T;
+ using High = T;
+ using Low = T;
+};
+
+
+//=================================================================================================
+//
+// MATHTRAIT SPECIALIZATION MACRO
+//
+//=================================================================================================
+
+//*************************************************************************************************
+/*! \cond internal */
+/*!\brief Macro for the creation of MathTrait specializations for the built-in data types.
+ * \ingroup math
+ *
+ * This macro is used for the setup of the MathTrait specializations for the built-in data
+ * types.
+ */
+#define PAIRS_CREATE_MATHTRAIT_SPECIALIZATION(T1,T2,HIGH,LOW) \
+ template<> \
+ struct MathTrait< T1, T2 > \
+ { \
+ typedef HIGH HighType; \
+ typedef LOW LowType; \
+ typedef HIGH High; \
+ typedef LOW Low; \
+ typedef HIGH AddType; \
+ typedef HIGH SubType; \
+ typedef HIGH MultType; \
+ typedef HIGH DivType; \
+ }
+/*! \endcond */
+//*************************************************************************************************
+
+
+
+
+//=================================================================================================
+//
+// UNSIGNED CHAR SPECIALIZATIONS
+//
+//=================================================================================================
+
+//*************************************************************************************************
+/*! \cond internal */
+// Type 1 Type 2 High type Low type
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned char , unsigned char , unsigned char , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned char , char , char , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned char , signed char , signed char , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned char , wchar_t , wchar_t , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned char , unsigned short, unsigned short, unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned char , short , short , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned char , unsigned int , unsigned int , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned char , int , int , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned char , unsigned long , unsigned long , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned char , long , long , unsigned char );
+#if defined(_WIN64)
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned char , std::size_t , std::size_t , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned char , std::ptrdiff_t, std::ptrdiff_t, unsigned char );
+#endif
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned char , float , float , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned char , double , double , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned char , long double , long double , unsigned char );
+/*! \endcond */
+//*************************************************************************************************
+
+
+
+
+//=================================================================================================
+//
+// CHAR SPECIALIZATIONS
+//
+//=================================================================================================
+
+//*************************************************************************************************
+/*! \cond internal */
+// Type 1 Type 2 High type Low type
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( char , unsigned char , char , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( char , char , char , char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( char , signed char , signed char , char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( char , wchar_t , wchar_t , char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( char , unsigned short, unsigned short, char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( char , short , short , char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( char , unsigned int , unsigned int , char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( char , int , int , char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( char , unsigned long , unsigned long , char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( char , long , long , char );
+#if defined(_WIN64)
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( char , std::size_t , std::size_t , char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( char , std::ptrdiff_t, std::ptrdiff_t, char );
+#endif
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( char , float , float , char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( char , double , double , char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( char , long double , long double , char );
+/*! \endcond */
+//*************************************************************************************************
+
+
+
+
+//=================================================================================================
+//
+// SIGNED CHAR SPECIALIZATIONS
+//
+//=================================================================================================
+
+//*************************************************************************************************
+/*! \cond internal */
+// Type 1 Type 2 High type Low type
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( signed char , unsigned char , signed char , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( signed char , char , signed char , char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( signed char , signed char , signed char , signed char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( signed char , wchar_t , wchar_t , signed char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( signed char , unsigned short, unsigned short, signed char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( signed char , short , short , signed char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( signed char , unsigned int , unsigned int , signed char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( signed char , int , int , signed char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( signed char , unsigned long , unsigned long , signed char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( signed char , long , long , signed char );
+#if defined(_WIN64)
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( signed char , std::size_t , std::size_t , signed char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( signed char , std::ptrdiff_t, std::ptrdiff_t, signed char );
+#endif
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( signed char , float , float , signed char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( signed char , double , double , signed char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( signed char , long double , long double , signed char );
+/*! \endcond */
+//*************************************************************************************************
+
+
+
+
+//=================================================================================================
+//
+// WCHAR_T SPECIALIZATIONS
+//
+//=================================================================================================
+
+//*************************************************************************************************
+/*! \cond internal */
+// Type 1 Type 2 High type Low type
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( wchar_t , unsigned char , wchar_t , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( wchar_t , char , wchar_t , char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( wchar_t , signed char , wchar_t , signed char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( wchar_t , wchar_t , wchar_t , wchar_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( wchar_t , unsigned short, unsigned short, wchar_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( wchar_t , short , short , wchar_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( wchar_t , unsigned int , unsigned int , wchar_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( wchar_t , int , int , wchar_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( wchar_t , unsigned long , unsigned long , wchar_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( wchar_t , long , long , wchar_t );
+#if defined(_WIN64)
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( wchar_t , std::size_t , std::size_t , wchar_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( wchar_t , std::ptrdiff_t, std::ptrdiff_t, wchar_t );
+#endif
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( wchar_t , float , float , wchar_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( wchar_t , double , double , wchar_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( wchar_t , long double , long double , wchar_t );
+/*! \endcond */
+//*************************************************************************************************
+
+
+
+
+//=================================================================================================
+//
+// UNSIGNED SHORT SPECIALIZATIONS
+//
+//=================================================================================================
+
+//*************************************************************************************************
+/*! \cond internal */
+// Type 1 Type 2 High type Low type
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned short, unsigned char , unsigned short, unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned short, char , unsigned short, char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned short, signed char , unsigned short, signed char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned short, wchar_t , unsigned short, wchar_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned short, unsigned short, unsigned short, unsigned short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned short, short , short , unsigned short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned short, unsigned int , unsigned int , unsigned short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned short, int , int , unsigned short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned short, unsigned long , unsigned long , unsigned short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned short, long , long , unsigned short );
+#if defined(_WIN64)
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned short, std::size_t , std::size_t , unsigned short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned short, std::ptrdiff_t, std::ptrdiff_t, unsigned short );
+#endif
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned short, float , float , unsigned short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned short, double , double , unsigned short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned short, long double , long double , unsigned short );
+/*! \endcond */
+//*************************************************************************************************
+
+
+
+
+//=================================================================================================
+//
+// SHORT SPECIALIZATIONS
+//
+//=================================================================================================
+
+//*************************************************************************************************
+/*! \cond internal */
+// Type 1 Type 2 High type Low type
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( short , unsigned char , short , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( short , char , short , char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( short , signed char , short , signed char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( short , wchar_t , short , wchar_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( short , unsigned short, short , unsigned short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( short , short , short , short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( short , unsigned int , unsigned int , short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( short , int , int , short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( short , unsigned long , unsigned long , short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( short , long , long , short );
+#if defined(_WIN64)
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( short , std::size_t , std::size_t , short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( short , std::ptrdiff_t, std::ptrdiff_t, short );
+#endif
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( short , float , float , short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( short , double , double , short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( short , long double , long double , short );
+/*! \endcond */
+//*************************************************************************************************
+
+
+
+
+//=================================================================================================
+//
+// UNSIGNED INT SPECIALIZATIONS
+//
+//=================================================================================================
+
+//*************************************************************************************************
+/*! \cond internal */
+// Type 1 Type 2 High type Low type
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned int , unsigned char , unsigned int , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned int , char , unsigned int , char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned int , signed char , unsigned int , signed char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned int , wchar_t , unsigned int , wchar_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned int , unsigned short, unsigned int , unsigned short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned int , short , unsigned int , short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned int , unsigned int , unsigned int , unsigned int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned int , int , int , unsigned int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned int , unsigned long , unsigned long , unsigned int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned int , long , long , unsigned int );
+#if defined(_WIN64)
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned int , std::size_t , std::size_t , unsigned int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned int , std::ptrdiff_t, std::ptrdiff_t, unsigned int );
+#endif
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned int , float , float , unsigned int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned int , double , double , unsigned int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned int , long double , long double , unsigned int );
+/*! \endcond */
+//*************************************************************************************************
+
+
+
+
+//=================================================================================================
+//
+// INT SPECIALIZATIONS
+//
+//=================================================================================================
+
+//*************************************************************************************************
+/*! \cond internal */
+// Type 1 Type 2 High type Low type
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( int , unsigned char , int , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( int , char , int , wchar_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( int , signed char , int , signed char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( int , wchar_t , int , wchar_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( int , unsigned short, int , unsigned short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( int , short , int , short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( int , unsigned int , int , unsigned int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( int , int , int , int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( int , unsigned long , unsigned long , int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( int , long , long , int );
+#if defined(_WIN64)
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( int , std::size_t , std::size_t , int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( int , std::ptrdiff_t, std::ptrdiff_t, int );
+#endif
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( int , float , float , int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( int , double , double , int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( int , long double , long double , int );
+/*! \endcond */
+//*************************************************************************************************
+
+
+
+
+//=================================================================================================
+//
+// UNSIGNED LONG SPECIALIZATIONS
+//
+//=================================================================================================
+
+//*************************************************************************************************
+/*! \cond internal */
+// Type 1 Type 2 High type Low type
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned long , unsigned char , unsigned long , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned long , char , unsigned long , char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned long , signed char , unsigned long , signed char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned long , wchar_t , unsigned long , wchar_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned long , unsigned short, unsigned long , unsigned short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned long , short , unsigned long , short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned long , unsigned int , unsigned long , unsigned int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned long , int , unsigned long , int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned long , unsigned long , unsigned long , unsigned long );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned long , long , long , unsigned long );
+#if defined(_WIN64)
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned long , std::size_t , std::size_t , unsigned long );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned long , std::ptrdiff_t, std::ptrdiff_t, unsigned long );
+#endif
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned long , float , float , unsigned long );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned long , double , double , unsigned long );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( unsigned long , long double , long double , unsigned long );
+/*! \endcond */
+//*************************************************************************************************
+
+
+
+
+//=================================================================================================
+//
+// LONG SPECIALIZATIONS
+//
+//=================================================================================================
+
+//*************************************************************************************************
+/*! \cond internal */
+// Type 1 Type 2 High type Low type
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long , unsigned char , long , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long , char , long , char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long , signed char , long , signed char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long , wchar_t , long , wchar_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long , unsigned short, long , unsigned short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long , short , long , short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long , unsigned int , long , unsigned int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long , int , long , int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long , unsigned long , long , unsigned long );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long , long , long , long );
+#if defined(_WIN64)
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long , std::size_t , std::size_t , long );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long , std::ptrdiff_t, std::ptrdiff_t, long );
+#endif
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long , float , float , long );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long , double , double , long );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long , long double , long double , long );
+/*! \endcond */
+//*************************************************************************************************
+
+
+
+
+//=================================================================================================
+//
+// SIZE_T SPECIALIZATIONS
+//
+//=================================================================================================
+
+//*************************************************************************************************
+#if defined(_WIN64)
+/*! \cond internal */
+// Type 1 Type 2 High type Low type
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( std::size_t , unsigned char , std::size_t , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( std::size_t , char , std::size_t , char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( std::size_t , signed char , std::size_t , signed char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( std::size_t , wchar_t , std::size_t , wchar_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( std::size_t , unsigned short, std::size_t , unsigned short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( std::size_t , short , std::size_t , short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( std::size_t , unsigned int , std::size_t , unsigned int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( std::size_t , int , std::size_t , int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( std::size_t , unsigned long , std::size_t , unsigned long );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( std::size_t , long , std::size_t , long );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( std::size_t , std::size_t , std::size_t , std::size_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( std::size_t , float , float , std::size_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( std::size_t , double , double , std::size_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( std::size_t , long double , long double , std::size_t );
+/*! \endcond */
+#endif
+//*************************************************************************************************
+
+
+
+
+//=================================================================================================
+//
+// FLOAT SPECIALIZATIONS
+//
+//=================================================================================================
+
+//*************************************************************************************************
+/*! \cond internal */
+// Type 1 Type 2 High type Low type
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( float , unsigned char , float , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( float , char , float , char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( float , signed char , float , signed char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( float , wchar_t , float , wchar_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( float , unsigned short, float , unsigned short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( float , short , float , short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( float , unsigned int , float , unsigned int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( float , int , float , int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( float , unsigned long , float , unsigned long );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( float , long , float , long );
+#if defined(_WIN64)
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( float , std::size_t , float , std::size_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( float , std::ptrdiff_t, float , std::ptrdiff_t );
+#endif
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( float , float , float , float );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( float , double , double , float );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( float , long double , long double , float );
+/*! \endcond */
+//*************************************************************************************************
+
+
+
+
+//=================================================================================================
+//
+// DOUBLE SPECIALIZATIONS
+//
+//=================================================================================================
+
+//*************************************************************************************************
+/*! \cond internal */
+// Type 1 Type 2 High type Low type
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( double , unsigned char , double , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( double , char , double , char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( double , signed char , double , signed char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( double , wchar_t , double , wchar_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( double , unsigned short, double , unsigned short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( double , short , double , short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( double , unsigned int , double , unsigned int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( double , int , double , int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( double , unsigned long , double , unsigned long );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( double , long , double , long );
+#if defined(_WIN64)
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( double , std::size_t , double , std::size_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( double , std::ptrdiff_t, double , std::ptrdiff_t );
+#endif
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( double , float , double , float );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( double , double , double , double );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( double , long double , long double , double );
+/*! \endcond */
+//*************************************************************************************************
+
+
+
+
+//=================================================================================================
+//
+// LONG DOUBLE SPECIALIZATIONS
+//
+//=================================================================================================
+
+//*************************************************************************************************
+/*! \cond internal */
+// Type 1 Type 2 High type Low type
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long double , unsigned char , long double , unsigned char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long double , char , long double , char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long double , signed char , long double , signed char );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long double , wchar_t , long double , wchar_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long double , unsigned short, long double , unsigned short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long double , short , long double , short );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long double , unsigned int , long double , unsigned int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long double , int , long double , int );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long double , unsigned long , long double , unsigned long );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long double , long , long double , long );
+#if defined(_WIN64)
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long double , std::size_t , long double , std::size_t );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long double , std::ptrdiff_t, long double , std::ptrdiff_t );
+#endif
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long double , float , long double , float );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long double , double , long double , double );
+PAIRS_CREATE_MATHTRAIT_SPECIALIZATION( long double , long double , long double , long double );
+/*! \endcond */
+//*************************************************************************************************
+
+#undef PAIRS_CREATE_MATHTRAIT_SPECIALIZATION
+
+}
diff --git a/runtime/math/Vector3.hpp b/runtime/math/Vector3.hpp
new file mode 100644
index 0000000..2e5b4f0
--- /dev/null
+++ b/runtime/math/Vector3.hpp
@@ -0,0 +1,43 @@
+#pragma once
+#include <iostream>
+
+#include "../pairs_common.hpp"
+#include "MathTrait.h"
+
+
+namespace pairs {
+
+#define HIGH typename MathTrait<Type,Other>::High
+
+template< typename Type >
+class Vector3 {
+public:
+ Vector3() = default;
+
+ // If the constructor is called from device, v_ is automatically allocated on
+ // device because it's a static array embeded in the object itself
+ PAIRS_ATTR_HOST_DEVICE Vector3( Type x, Type y, Type z ) {
+ v_[0] = x;
+ v_[1] = y;
+ v_[2] = z;
+ }
+
+ template< typename Other >
+ PAIRS_ATTR_HOST_DEVICE inline Vector3<HIGH> operator+( const Vector3<Other>& rhs ) const{
+ return Vector3<HIGH>( v_[0]+static_cast<Type>(rhs.v_[0]), v_[1]+static_cast<Type>(rhs.v_[1]), v_[2]+static_cast<Type>(rhs.v_[2]) );
+ }
+
+ PAIRS_ATTR_HOST_DEVICE Type& operator[]( int index ) {
+ return v_[index];
+ }
+
+ PAIRS_ATTR_HOST_DEVICE const Type& operator[] ( int index ) const {
+ return v_[index];
+ }
+
+private:
+ Type v_[3] = {Type(), Type(), Type()};
+};
+#undef HIGH
+
+}
diff --git a/runtime/pairs.hpp b/runtime/pairs.hpp
index 27892e4..cf1e65e 100644
--- a/runtime/pairs.hpp
+++ b/runtime/pairs.hpp
@@ -239,6 +239,10 @@ public:
void copyPropertyToHost(Property &prop, action_t action, size_t size);
+ DeviceFlags* getPropFlags(){
+ return prop_flags;
+ }
+
// Contact properties
ContactProperty &getContactProperty(property_t id);
ContactProperty &getContactPropertyByName(std::string name);
diff --git a/runtime/pairs_common.hpp b/runtime/pairs_common.hpp
index 23a03ba..73c9f7b 100644
--- a/runtime/pairs_common.hpp
+++ b/runtime/pairs_common.hpp
@@ -5,6 +5,16 @@
namespace pairs {
+#ifdef PAIRS_TARGET_CUDA
+ #define PAIRS_ATTR_HOST __host__
+ #define PAIRS_ATTR_DEVICE __device__
+ #define PAIRS_ATTR_HOST_DEVICE __host__ __device__
+#else
+ #define PAIRS_ATTR_HOST
+ #define PAIRS_ATTR_DEVICE
+ #define PAIRS_ATTR_HOST_DEVICE
+#endif
+
constexpr int FLAGS_INFINITE = 1 << 0 ;
constexpr int FLAGS_GHOST = 1 << 1 ;
constexpr int FLAGS_FIXED = 1 << 2 ;
diff --git a/src/pairs/code_gen/accessor.py b/src/pairs/code_gen/accessor.py
new file mode 100644
index 0000000..9996a6f
--- /dev/null
+++ b/src/pairs/code_gen/accessor.py
@@ -0,0 +1,369 @@
+from pairs.ir.types import Types
+
+class PairsAcessor:
+ def __init__(self, cgen):
+ self.sim = cgen.sim
+ self.target = cgen.target
+ self.print = cgen.print
+ self.debug = cgen.debug
+ self.host_device_attr = ""
+ self.host_attr = ""
+
+ def generate(self):
+ self.print("")
+
+ if self.target.is_gpu():
+ self.host_device_attr = "__host__ __device__ "
+ self.host_attr = "__host__ "
+ self.print("#include \"runtime/math/Vector3.hpp\"")
+ # self.print("#include \"runtime/math/Quaternion.hpp\"")
+ # self.print("#include \"runtime/math/Matrix3.hpp\"")
+ self.print("")
+
+ self.print("class PairsAccessor {")
+ self.print("private:")
+ self.print.add_indent(4)
+ self.member_variables()
+ self.print.add_indent(-4)
+ self.print("public:")
+ self.print.add_indent(4)
+
+ if self.target.is_gpu():
+ self.update()
+
+ self.constructor()
+
+ for p in self.sim.properties:
+ if (p.type()==Types.Vector) or (Types.is_scalar(p.type())):
+ self.get_property(p)
+ self.set_property(p)
+ if self.target.is_gpu():
+ self.sync_property(p)
+
+ self.utility_funcs()
+
+ self.print.add_indent(-4)
+ self.print("};")
+ self.print("")
+
+ def member_variables(self):
+ self.print("PairsSimulation *ps;")
+
+ if self.target.is_gpu():
+ self.print("int *nlocal_d;")
+ self.print("int *nghost_d;")
+ self.print("")
+
+ self.print("//Properties")
+ for p in self.sim.properties:
+ pname = p.name()
+ tkw = Types.c_keyword(self.sim, p.type())
+ self.print(f"{tkw} *{pname}_d;")
+
+ self.print("")
+ self.print("//Property flags")
+ for p in self.sim.properties:
+ pname = p.name()
+ tkw = Types.c_keyword(self.sim, Types.Boolean)
+ self.print(f"{tkw} *{pname}_device_flag_d;")
+ self.print(f"{tkw} {pname}_device_flag_h = false;")
+ self.print(f"{tkw} {pname}_host_flag = false;")
+
+ self.print("")
+
+ def update(self):
+ self.print(f"{self.host_attr}void update(){{")
+ self.print.add_indent(4)
+ self.print(f"cudaMemcpy(nlocal_d, &(ps->pobj->nlocal), sizeof(int), cudaMemcpyHostToDevice);")
+ self.print(f"cudaMemcpy(nghost_d, &(ps->pobj->nghost), sizeof(int), cudaMemcpyHostToDevice);")
+
+ for p in self.sim.properties:
+ pname = p.name()
+ self.print(f"{pname}_d = ps->pobj->{pname}_d;")
+
+ self.print.add_indent(-4)
+ self.print("}")
+ self.print("")
+
+
+ def constructor(self):
+ if self.target.is_gpu():
+ self.print(f"{self.host_attr}PairsAccessor(PairsSimulation *ps_): ps(ps_){{")
+ self.print.add_indent(4)
+
+ self.print(f"cudaMalloc(&nlocal_d, sizeof(int));")
+ self.print(f"cudaMalloc(&nghost_d, sizeof(int));")
+ self.print("this->update();")
+
+ for p in self.sim.properties:
+ pname = p.name()
+ tkw = Types.c_keyword(self.sim, Types.Boolean)
+ self.print(f"cudaMalloc(&{pname}_device_flag_d, sizeof({tkw}));")
+ self.print(f"cudaMemcpy({pname}_device_flag_d, &{pname}_device_flag_h, sizeof({tkw}), cudaMemcpyHostToDevice);")
+
+ self.print.add_indent(-4)
+ self.print("}")
+ else:
+ self.print("PairsAccessor(PairsSimulation *ps_): ps(ps_){}")
+
+ self.print("")
+
+ def ifdef_else(self, ifdef, func1, args1, func2, args2):
+ self.print.add_indent(4)
+ self.print(f"#ifdef {ifdef}")
+ func1(*args1)
+ self.print("#else")
+ func2(*args2)
+ self.print("#endif")
+ self.print.add_indent(-4)
+
+ def getter_body(self, prop, device=False):
+ self.print.add_indent(4)
+ pname = prop.name()
+ tkw = Types.c_accessor_keyword(self.sim, prop.type())
+
+ if self.target.is_gpu() and device:
+ v = f"{pname}_d"
+ else:
+ v = f"ps->pobj->{pname}"
+
+ if Types.is_scalar(prop.type()):
+ self.print(f"return {v}[i];")
+ else:
+ nelems = Types.number_of_elements(self.sim, prop.type())
+ return_values = [f"{v}[i*{nelems} + {n}]" for n in range(nelems)]
+ self.print(f"return {tkw}(" + ", ".join(rv for rv in return_values) + ");")
+ self.print.add_indent(-4)
+
+
+ def get_property(self, prop):
+ pname = prop.name()
+ tkw = Types.c_accessor_keyword(self.sim, prop.type())
+ splitname = pname.split('_')
+ funcname = ''.join(word.capitalize() for word in splitname)
+
+ self.print(f"{self.host_device_attr}{tkw} get{funcname}(const size_t i) const{{")
+
+ if self.target.is_gpu():
+ self.ifdef_else("__CUDA_ARCH__", self.getter_body, [prop, True], self.getter_body, [prop, False])
+ else:
+ self.getter_body(prop, False)
+
+ self.print("}")
+ self.print("")
+
+
+ def setter_body(self, prop, device=False):
+ self.print.add_indent(4)
+ pname = prop.name()
+ tkw = Types.c_accessor_keyword(self.sim, prop.type())
+
+ if self.target.is_gpu() and device:
+ v = f"{pname}_d"
+ else:
+ v = f"ps->pobj->{pname}"
+
+ if Types.is_scalar(prop.type()):
+ self.print(f"{v}[i] = value;")
+ else:
+ nelems = Types.number_of_elements(self.sim, prop.type())
+ set_values = [f"{v}[i*{nelems} + {n}] = value[{n}];" for n in range(nelems)]
+ for sv in set_values:
+ self.print(sv)
+
+ if self.target.is_gpu():
+ flag = f"*{pname}_device_flag_d" if device else f"{pname}_host_flag"
+ self.print(f"{flag} = true;")
+
+ self.print.add_indent(-4)
+
+
+ def set_property(self, prop):
+ pname = prop.name()
+ tkw = Types.c_accessor_keyword(self.sim, prop.type())
+ splitname = pname.split('_')
+ funcname = ''.join(word.capitalize() for word in splitname)
+
+ self.print(f"{self.host_device_attr}void set{funcname}(const size_t i, const {tkw} &value){{")
+
+ if self.target.is_gpu():
+ self.ifdef_else("__CUDA_ARCH__", self.setter_body, [prop, True], self.setter_body, [prop, False])
+ else:
+ self.setter_body(prop, False)
+
+ self.print("}")
+ self.print("")
+
+ def sync_property(self, prop):
+ pname = prop.name()
+ pid = prop.id()
+ splitname = pname.split('_')
+ funcname = ''.join(word.capitalize() for word in splitname)
+
+ self.print(f"{self.host_attr}void sync{funcname}(){{")
+ self.print.add_indent(4)
+ self.print(f"{pname}_d = ps->pobj->{pname}_d;")
+ self.print(f"cudaMemcpy(&{pname}_device_flag_h, {pname}_device_flag_d, sizeof(bool), cudaMemcpyDeviceToHost);")
+ self.print("")
+
+
+ #####################################################################################################################
+ #####################################################################################################################
+ # self.print(f"if (({pname}_host_flag && {pname}_device_flag_h) || ")
+ # self.print.add_indent(4)
+ # self.print(f"({pname}_host_flag && !ps->pairs_runtime->getPropFlags()->isHostFlagSet({pid})) ||")
+ # self.print(f"({pname}_device_flag_h && !ps->pairs_runtime->getPropFlags()->isDeviceFlagSet({pid}))){{")
+ # self.print(f"PAIRS_ERROR(\"OUT OF SYNC! Both host and device versions of {pname} are in a modified state.\\n\");")
+ # self.print("exit(-1);")
+ # self.print.add_indent(-4)
+ # self.print("}")
+ # self.print("")
+
+
+ self.print(f"if ({pname}_host_flag && {pname}_device_flag_h){{")
+ self.print.add_indent(4)
+ self.print(f"PAIRS_ERROR(\"OUT OF SYNC 1! Both host and device versions of {pname} are in a modified state.\\n\");")
+ self.print("exit(-1);")
+ self.print.add_indent(-4)
+ self.print("}")
+ self.print("")
+
+ self.print(f"if ({pname}_host_flag && !ps->pairs_runtime->getPropFlags()->isHostFlagSet({pid})){{")
+ self.print.add_indent(4)
+ self.print(f"PAIRS_ERROR(\"OUT OF SYNC 2! Both host and device versions of {pname} are in a modified state.\\n\");")
+ self.print("exit(-1);")
+ self.print.add_indent(-4)
+ self.print("}")
+ self.print("")
+
+ self.print(f"if ({pname}_device_flag_h && !ps->pairs_runtime->getPropFlags()->isDeviceFlagSet({pid})){{")
+ self.print.add_indent(4)
+ self.print(f"PAIRS_ERROR(\"OUT OF SYNC 3! Both host and device versions of {pname} are in a modified state.\\n\");")
+ self.print("exit(-1);")
+ self.print.add_indent(-4)
+ self.print("}")
+ self.print("")
+
+ #####################################################################################################################
+ #####################################################################################################################
+
+
+ self.print(f"if ({pname}_host_flag){{")
+ self.print.add_indent(4)
+ self.print(f"ps->pairs_runtime->getPropFlags()->setHostFlag({pid});")
+ self.print(f"ps->pairs_runtime->getPropFlags()->clearDeviceFlag({pid});")
+ self.print.add_indent(-4)
+ self.print("}")
+
+ self.print(f"else if ({pname}_device_flag_h){{")
+ self.print.add_indent(4)
+ self.print(f"ps->pairs_runtime->getPropFlags()->setDeviceFlag({pid});")
+ self.print(f"ps->pairs_runtime->getPropFlags()->clearHostFlag({pid});")
+ self.print.add_indent(-4)
+ self.print("}")
+ self.print("")
+
+ nelems = Types.number_of_elements(self.sim, prop.type())
+ tkw = Types.c_keyword(self.sim, prop.type())
+
+ self.print(f"if (ps->pairs_runtime->getPropFlags()->isHostFlagSet({pid})) {{")
+ self.print.add_indent(4)
+
+ self.print(f"ps->pairs_runtime->copyPropertyToDevice({pid}, ReadOnly, (((ps->pobj->nlocal + ps->pobj->nghost) * {nelems}) * sizeof({tkw})));")
+ self.print.add_indent(-4)
+ self.print("}")
+
+ self.print(f"else if (ps->pairs_runtime->getPropFlags()->isDeviceFlagSet({pid})) {{")
+ self.print.add_indent(4)
+ self.print(f"ps->pairs_runtime->copyPropertyToHost({pid}, ReadOnly, (((ps->pobj->nlocal + ps->pobj->nghost) * {nelems}) * sizeof({tkw})));")
+ self.print.add_indent(-4)
+ self.print("}")
+ self.print("")
+
+ self.print(f"{pname}_host_flag = false;")
+ self.print(f"{pname}_device_flag_h = false;")
+ self.print(f"cudaMemcpy({pname}_device_flag_d, &{pname}_device_flag_h, sizeof(bool), cudaMemcpyHostToDevice);")
+
+ self.print.add_indent(-4)
+ self.print("}")
+ self.print("")
+
+ def utility_funcs(self):
+ if self.target.is_gpu():
+ self.print(f"{self.host_device_attr}int size() const {{")
+ self.print(" #ifdef __CUDA_ARCH__")
+ self.print(" return *nlocal_d + *nghost_d;")
+ self.print(" #else")
+ self.print(" return ps->pobj->nlocal + ps->pobj->nghost;")
+ self.print(" #endif")
+ self.print("}")
+ self.print("")
+ else:
+ self.print("int size() const {return ps->pobj->nlocal + ps->pobj->nghost;}")
+
+ if self.target.is_gpu():
+ self.print(f"{self.host_device_attr}int nlocal() const {{")
+ self.print(" #ifdef __CUDA_ARCH__")
+ self.print(" return *nlocal_d;")
+ self.print(" #else")
+ self.print(" return ps->pobj->nlocal;")
+ self.print(" #endif")
+ self.print("}")
+ self.print("")
+ else:
+ self.print("int nlocal() const {return ps->pobj->nlocal;}")
+
+ if self.target.is_gpu():
+ self.print(f"{self.host_device_attr}int nghost() const {{")
+ self.print(" #ifdef __CUDA_ARCH__")
+ self.print(" return *nghost_d;")
+ self.print(" #else")
+ self.print(" return ps->pobj->nghost;")
+ self.print(" #endif")
+ self.print("}")
+ self.print("")
+ else:
+ self.print("int nghost() const {return ps->pobj->nghost;}")
+
+
+ self.print(f"{self.host_device_attr}int getInvalidIdx(){{return -1;}}")
+ self.print("")
+
+ self.print(f"{self.host_device_attr}pairs::id_t getInvalidUid(){{return 0;}}")
+ self.print("")
+
+ self.print(f"{self.host_device_attr}int uidToIdx(pairs::id_t uid){{")
+ self.print(" int idx = getInvalidIdx();")
+ self.print(" for(int i=0; i<size(); ++i){")
+ self.print(" if (getUid(i) == uid){")
+ self.print(" idx = i;")
+ self.print(" break;")
+ self.print(" }")
+ self.print(" }")
+ self.print(" return idx;")
+ self.print("}")
+ self.print("")
+
+ self.print(f"{self.host_device_attr}int uidToIdxLocal(pairs::id_t uid){{")
+ self.print(" int idx = getInvalidIdx();")
+ self.print(" for(int i=0; i<nlocal(); ++i){")
+ self.print(" if (getUid(i) == uid){")
+ self.print(" idx = i;")
+ self.print(" break;")
+ self.print(" }")
+ self.print(" }")
+ self.print(" return idx;")
+ self.print("}")
+ self.print("")
+
+ self.print(f"{self.host_device_attr}int uidToIdxGhost(pairs::id_t uid){{")
+ self.print(" int idx = getInvalidIdx();")
+ self.print(" for(int i=nlocal(); i<size(); ++i){")
+ self.print(" if (getUid(i) == uid){")
+ self.print(" idx = i;")
+ self.print(" break;")
+ self.print(" }")
+ self.print(" }")
+ self.print(" return idx;")
+ self.print("}")
+ self.print("")
diff --git a/src/pairs/code_gen/cgen.py b/src/pairs/code_gen/cgen.py
index 766dcba..29e8b2f 100644
--- a/src/pairs/code_gen/cgen.py
+++ b/src/pairs/code_gen/cgen.py
@@ -33,6 +33,7 @@ from pairs.ir.vectors import Vector, VectorAccess, VectorOp, ZeroVector
from pairs.sim.domain_partitioners import DomainPartitioners
from pairs.sim.timestep import Timestep
from pairs.code_gen.printer import Printer
+from pairs.code_gen.accessor import PairsAcessor
class CGen:
@@ -61,7 +62,7 @@ class CGen:
# Ideally this should never be called
return "nullptr"
- name = obj.name() if not device else f"d_{obj.name()}"
+ name = obj.name() if not device else f"{obj.name()}_d"
t = obj.type()
if not Types.is_scalar(t) and index is not None:
name += f"_{index}"
@@ -197,102 +198,8 @@ class CGen:
self.print.add_indent(-4)
self.print("}")
self.print("")
-
- def generate_host_pairs_accessor_class(self):
- self.print("#include <core/math/Vector3.h>")
- self.print("#include <core/math/Quaternion.h>")
- self.print("#include <core/math/Matrix3.h>")
- self.print("")
- self.print("class PairsAccessor {")
- self.print("private:")
- self.print(" std::shared_ptr<PairsSimulation> ps;")
- self.print("public:")
- self.print.add_indent(4)
- self.print("PairsAccessor(const std::shared_ptr<PairsSimulation> &ps_): ps(ps_){}")
- self.print("")
- self.print("int size() const { return ps->pobj->nlocal + ps->pobj->nghost; }")
- self.print("")
-
- self.print("int nlocal() const { return ps->pobj->nlocal; }")
- self.print("")
-
- self.print("int nghost() const { return ps->pobj->nghost; }")
- self.print("")
-
- self.print("int getInvalidIdx(){return -1;}")
- self.print("")
-
- self.print("pairs::id_t getInvalidUid(){return 0;}")
- self.print("")
-
- self.print('''int uidToIdx(pairs::id_t uid){
- int idx = getInvalidIdx();
- for(int i=0; i<this->nlocal(); ++i){
- if (getUid(i) == uid){
- idx = i;
- break;
- }
- }
- return idx;''')
- self.print("}")
- self.print("")
-
- self.print('''int uidToIdxGhost(pairs::id_t uid){
- int idx = getInvalidIdx();
- for(int i=this->nlocal(); i<this->size(); ++i){
- if (getUid(i) == uid){
- idx = i;
- break;
- }
- }
- return idx;''')
- self.print("}")
- self.print("")
-
- for p in self.sim.properties:
- ptr = p.name()
-
- tkw = Types.walberla_keyword(self.sim, p.type())
-
- splitname = ptr.split('_')
- funcname = ''.join(word.capitalize() for word in splitname)
- v = f"ps->pobj->{ptr}"
- getSignature = f"{tkw} get{funcname}(const size_t i) const"
- setSignature = f"void set{funcname}(const size_t i, {tkw} const &value)"
-
- if Types.is_scalar(p.type()):
- self.print(f"{getSignature} {{return {v}[i];}}")
- self.print(f"{setSignature} {{{v}[i] = value;}}")
- self.print("")
-
- elif p.type()==Types.Vector:
- self.print(f"{getSignature} {{return {tkw}({v}[i*3 + 0], {v}[i*3 + 1], {v}[i*3 + 2]);}}")
- self.print(f"{setSignature} {{{v}[i*3 + 0] = value[0]; {v}[i*3 + 1] = value[1]; {v}[i*3 + 2] = value[2];}}")
- self.print("")
-
- elif p.type()==Types.Quaternion:
- self.print(f"{getSignature} {{return {tkw}({v}[i*4 + 0], {v}[i*4 + 1], {v}[i*4 + 2], {v}[i*4 + 3]);}}")
- self.print(f"{setSignature} {{{v}[i*4 + 0] = value[0]; {v}[i*4 + 1] = value[1]; {v}[i*4 + 2] = value[2]; {v}[i*4 + 3] = value[3];}}")
- self.print("")
-
- elif p.type()==Types.Matrix:
- self.print(f"{getSignature} {{" \
- f"return {tkw}({v}[i*9 + 0], {v}[i*9 + 1], {v}[i*9 + 2], "\
- f"{v}[i*9 + 3], {v}[i*9 + 4], {v}[i*9 + 5], "\
- f"{v}[i*9 + 6], {v}[i*9 + 7], {v}[i*9 + 8]);}}")
- self.print(f"{setSignature} {{" \
- f"{v}[i*9 + 0] = value[0]; {v}[i*9 + 1] = value[1]; {v}[i*9 + 2] = value[2]; "\
- f"{v}[i*9 + 3] = value[3]; {v}[i*9 + 4] = value[4]; {v}[i*9 + 5] = value[5]; "\
- f"{v}[i*9 + 6] = value[6]; {v}[i*9 + 7] = value[7]; {v}[i*9 + 8] = value[8];}}")
- self.print("")
-
-
- self.print.add_indent(-4)
- self.print("};")
- self.print("")
-
+
def generate_pairs_object_structure(self):
-
self.print("")
externkw = "" if self.sim._generate_whole_program else "extern "
if self.target.is_gpu():
@@ -301,14 +208,14 @@ class CGen:
t = array.type()
tkw = Types.c_keyword(self.sim, t)
size = self.generate_expression(ScalarOp.inline(array.alloc_size()))
- self.print(f"{externkw}__constant__ {tkw} d_{array.name()}[{size}];")
+ self.print(f"{externkw}__constant__ {tkw} {array.name()}_d[{size}];")
for feature_prop in self.sim.feature_properties:
if feature_prop.device_flag:
t = feature_prop.type()
tkw = Types.c_keyword(self.sim, t)
size = feature_prop.array_size()
- self.print(f"{externkw}__constant__ {tkw} d_{feature_prop.name()}[{size}];")
+ self.print(f"{externkw}__constant__ {tkw} {feature_prop.name()}_d[{size}];")
self.print("")
self.print("struct PairsObjects {")
@@ -330,7 +237,7 @@ class CGen:
if a.is_static():
continue
else:
- self.print(f"{tkw} *d_{ptr};")
+ self.print(f"{tkw} *{ptr}_d;")
self.print("// Properties")
for p in self.sim.properties:
@@ -339,7 +246,7 @@ class CGen:
self.print(f"{tkw} *{ptr};")
if self.target.is_gpu() and p.device_flag:
- self.print(f"{tkw} *d_{ptr};")
+ self.print(f"{tkw} *{ptr}_d;")
self.print("// Contact properties")
for cp in self.sim.contact_properties:
@@ -348,7 +255,7 @@ class CGen:
self.print(f"{tkw} *{ptr};")
if self.target.is_gpu() and cp.device_flag:
- self.print(f"{tkw} *d_{ptr};")
+ self.print(f"{tkw} *{ptr}_d;")
self.print("// Feature properties")
for fp in self.sim.feature_properties:
@@ -423,14 +330,14 @@ class CGen:
t = array.type()
tkw = Types.c_keyword(self.sim, t)
size = self.generate_expression(ScalarOp.inline(array.alloc_size()))
- self.print(f"__constant__ {tkw} d_{array.name()}[{size}];")
+ self.print(f"__constant__ {tkw} {array.name()}_d[{size}];")
for feature_prop in self.sim.feature_properties:
if feature_prop.device_flag:
t = feature_prop.type()
tkw = Types.c_keyword(self.sim, t)
size = feature_prop.array_size()
- self.print(f"__constant__ {tkw} d_{feature_prop.name()}[{size}];")
+ self.print(f"__constant__ {tkw} {feature_prop.name()}_d[{size}];")
self.print("")
@@ -509,9 +416,16 @@ class CGen:
self.print("}")
self.print("")
- self.print("int rank(){")
- self.print(" return pairs_runtime->getDomainPartitioner()->getRank();")
- self.print("}")
+ self.print("int rank(){ return pairs_runtime->getDomainPartitioner()->getRank();}")
+ self.print("")
+
+ self.print("int size(){ return pobj->nlocal + pobj->nghost;}")
+ self.print("")
+
+ self.print("int nlocal(){ return pobj->nlocal;}")
+ self.print("")
+
+ self.print("int nghost(){ return pobj->nghost;}")
self.print("")
self.print("void setup_sim() {")
@@ -529,7 +443,7 @@ class CGen:
self.print("}")
self.print("")
- self.print("void communicate(int timestep) {")
+ self.print("void communicate(int timestep = 0) {")
self.print(" pobj->sim_timestep = timestep;")
self.generate_statement(communicate_module.block)
self.print("}")
@@ -555,9 +469,7 @@ class CGen:
self.print.add_indent(-4)
self.print("};")
- if self.sim.partitioner()==DomainPartitioners.BlockForest:
- self.generate_host_pairs_accessor_class()
- # self.generate_host_pairs_accessor_class()
+ PairsAcessor(self).generate()
self.print.end()
self.generate_full_object_names = False
@@ -581,33 +493,33 @@ class CGen:
for array in module.arrays():
type_kw = Types.c_keyword(self.sim, array.type())
- name = array.name() if not device_cond else f"d_{array.name()}"
+ name = array.name() if not device_cond else f"{array.name()}_d"
if not array.is_static() or (array.is_static() and not device_cond):
self.print(f"{type_kw} *{array.name()} = pobj->{name};")
if array in module.host_references():
- self.print(f"{type_kw} *h_{array.name()} = pobj->{array.name()};")
+ self.print(f"{type_kw} *{array.name()}_h = pobj->{array.name()};")
for prop in module.properties():
type_kw = Types.c_keyword(self.sim, prop.type())
- name = prop.name() if not device_cond else f"d_{prop.name()}"
+ name = prop.name() if not device_cond else f"{prop.name()}_d"
self.print(f"{type_kw} *{prop.name()} = pobj->{name};")
if prop in module.host_references():
- self.print(f"{type_kw} *h_{prop.name()} = pobj->{prop.name()};")
+ self.print(f"{type_kw} *{prop.name()}_h = pobj->{prop.name()};")
for contact_prop in module.contact_properties():
type_kw = Types.c_keyword(self.sim, contact_prop.type())
- name = contact_prop.name() if not device_cond else f"d_{contact_prop.name()}"
+ name = contact_prop.name() if not device_cond else f"{contact_prop.name()}_d"
self.print(f"{type_kw} *{contact_prop.name()} = pobj->{name};")
if contact_prop in module.host_references():
- self.print(f"{type_kw} *h_{contact_prop.name()} = pobj->{contact_prop.name()};")
+ self.print(f"{type_kw} *{contact_prop.name()}_h = pobj->{contact_prop.name()};")
for feature_prop in module.feature_properties():
type_kw = Types.c_keyword(self.sim, feature_prop.type())
- name = feature_prop.name() if not device_cond else f"d_{feature_prop.name()}"
+ name = feature_prop.name() if not device_cond else f"{feature_prop.name()}_d"
if feature_prop.device_flag and device_cond:
# self.print(f"{type_kw} *{feature_prop.name()} = {self.generate_object_reference(feature_prop, device=device_cond)};")
@@ -616,7 +528,7 @@ class CGen:
self.print(f"{type_kw} *{feature_prop.name()} = pobj->{name};")
if feature_prop in module.host_references():
- self.print(f"{type_kw} *h_{feature_prop.name()} = pobj->{feature_prop.name()};")
+ self.print(f"{type_kw} *{feature_prop.name()}_h = pobj->{feature_prop.name()};")
def generate_main(self, module):
assert module.name=='main'
@@ -1018,11 +930,11 @@ class CGen:
if ast_node.decl:
self.print(f"{tkw} *{array_name} = ({tkw} *) malloc({size});")
if self.target.is_gpu() and ast_node.array.device_flag:
- self.print(f"{tkw} *d_{array_name} = ({tkw} *) pairs::device_alloc({size});")
+ self.print(f"{tkw} *{array_name}_d = ({tkw} *) pairs::device_alloc({size});")
else:
self.print(f"{array_name} = ({tkw} *) malloc({size});")
if self.target.is_gpu() and ast_node.array.device_flag:
- self.print(f"d_{array_name} = ({tkw} *) pairs::device_alloc({size});")
+ self.print(f"{array_name}_d = ({tkw} *) pairs::device_alloc({size});")
if isinstance(ast_node, KernelLaunch):
range_start = self.generate_expression(ScalarOp.inline(ast_node.min))
@@ -1274,14 +1186,14 @@ class CGen:
if isinstance(ast_node, DeviceStaticRef):
elem = self.generate_expression(ast_node.elem)
- return f"d_{elem}"
+ return f"{elem}_d"
if isinstance(ast_node, FeatureProperty):
return self.generate_object_reference(ast_node)
if isinstance(ast_node, HostRef):
elem = self.generate_expression(ast_node.elem)
- return f"h_{elem}"
+ return f"{elem}_h"
if isinstance(ast_node, Iter):
assert mem is False, "Iterator is not lvalue!"
diff --git a/src/pairs/ir/types.py b/src/pairs/ir/types.py
index e1b3b51..f4ab048 100644
--- a/src/pairs/ir/types.py
+++ b/src/pairs/ir/types.py
@@ -13,13 +13,13 @@ class Types:
Matrix = 10
Quaternion = 11
- def walberla_keyword(sim, t):
+ def c_accessor_keyword(sim, t):
real_kw = 'double' if sim.use_double_precision() else 'float'
return (
real_kw if t==Types.Real
- else f'walberla::math::Vector3<{real_kw}>' if t==Types.Vector
- else f'walberla::math::Matrix3<{real_kw}>' if t==Types.Matrix
- else f'walberla::math::Quaternion<{real_kw}>' if t==Types.Quaternion
+ else f'pairs::Vector3<{real_kw}>' if t==Types.Vector
+ else f'pairs::Matrix3<{real_kw}>' if t==Types.Matrix
+ else f'pairs::Quaternion<{real_kw}>' if t==Types.Quaternion
else 'float' if t == Types.Float
else 'double' if t == Types.Double
else 'int' if t == Types.Int32
--
GitLab