From c818bbcc27d23c7ed12be2e4b334dc107388d303 Mon Sep 17 00:00:00 2001 From: Serhii Yaskovets <yaskovet@mpi-cbg.de> Date: Fri, 20 Oct 2023 14:39:19 +0200 Subject: [PATCH] Add support for 3-index vector in vector_dist_exression --- .../cuda/vector_dist_operators_cuda.cuh | 81 +++++++++++++++++++ .../Vector/vector_dist_operators.hpp | 4 +- 2 files changed, 83 insertions(+), 2 deletions(-) diff --git a/src/Operators/Vector/cuda/vector_dist_operators_cuda.cuh b/src/Operators/Vector/cuda/vector_dist_operators_cuda.cuh index 28cac20..d05b322 100644 --- a/src/Operators/Vector/cuda/vector_dist_operators_cuda.cuh +++ b/src/Operators/Vector/cuda/vector_dist_operators_cuda.cuh @@ -280,6 +280,47 @@ struct get_vector_dist_expression_op<1,false> } }; +template<> +struct get_vector_dist_expression_op<2,false> +{ + template<typename exp_type> + __device__ __host__ static int get(exp_type & o1, const vect_dist_key_dx & key, const int (& comp)[2]) + { + printf("ERROR: Slicer, the expression is incorrect, please check it\n"); + return 0; + } + + template<unsigned int prop, typename exp_type, typename vector_type> + __device__ __host__ inline static void assign(exp_type & o1, vector_type & v, const vect_dist_key_dx & key, const vect_dist_key_dx & key_orig, const int (& comp)[2]) + { + printf("ERROR: Slicer, the expression is incorrect, please check it\n"); + } + + template<unsigned int prop, typename exp_type, typename vector_type> + __device__ __host__ inline static void assign(exp_type & o1, vector_type & v, const unsigned int & key, const vect_dist_key_dx & key_orig, const int (& comp)[2]) + { + printf("ERROR: Slicer, the expression is incorrect, please check it\n"); + } + + template<unsigned int prop,typename exp_type, typename vector_type> + __device__ __host__ inline static void assign(exp_type & o1, vector_type & v, const vect_dist_key_dx & key, const vect_dist_key_dx & key_orig, const Point<2,int> & comp) + { + printf("ERROR: Slicer, the expression is incorrect, please check it\n"); + } + + template<unsigned int prop,typename exp_type, typename vector_type> + __device__ __host__ inline static void assign(exp_type & o1, vector_type & v, const unsigned int & key, const unsigned int & key_orig, const Point<2,int> & comp) + { + printf("ERROR: Slicer, the expression is incorrect, please check it\n"); + } + + template<unsigned int prop, typename vector_type> + inline static void assign_double(double d, vector_type & v, const vect_dist_key_dx & key, const int (& comp)[2]) + { + printf("ERROR: Slicer, the expression is incorrect, please check it\n"); + } +}; + template<> struct get_vector_dist_expression_op<1,true> { @@ -360,6 +401,46 @@ struct get_vector_dist_expression_op<2,true> } }; +template<> +struct get_vector_dist_expression_op<3,true> +{ + template<typename exp_type> + __device__ __host__ static auto get(exp_type & o1, const vect_dist_key_dx & key, const int (& comp)[3]) -> decltype(o1.value(vect_dist_key_dx(0))[0][0][0]) + { + return o1.value(key)[comp[0]][comp[1]][comp[2]]; + } + + template<unsigned int prop,typename exp_type, typename vector_type> + __device__ __host__ inline static void assign(exp_type & o1, vector_type & v, const vect_dist_key_dx & key, const vect_dist_key_dx & key_orig, const int (& comp)[3]) + { + pos_or_propL<vector_type,prop>::value(v,key)[comp[0]][comp[1]][comp[2]] = o1.value(key_orig); + } + + template<unsigned int prop,typename exp_type, typename vector_type> + __device__ __host__ inline static void assign(exp_type & o1, vector_type & v, const vect_dist_key_dx & key, const vect_dist_key_dx & key_orig, const Point<3,int> & comp) + { + pos_or_propL<vector_type,prop>::value(v,key)[comp[0]][comp[1]][comp[2]] = o1.value(key_orig); + } + + template<unsigned int prop,typename exp_type, typename vector_type> + __device__ __host__ inline static void assign(exp_type & o1, vector_type & v, const unsigned int & key, const unsigned int & key_orig, const int (& comp)[3]) + { + pos_or_propL<vector_type,prop>::value(v,key)[comp[0]][comp[1]][comp[2]] = o1.value(key_orig); + } + + template<unsigned int prop,typename exp_type, typename vector_type> + __device__ __host__ inline static void assign(exp_type & o1, vector_type & v, const unsigned int & key, const unsigned int & key_orig, const Point<3,int> & comp) + { + pos_or_propL<vector_type,prop>::value(v,key)[comp[0]][comp[1]][comp[2]] = o1.value(key_orig); + } + + template<unsigned int prop, typename vector_type> + inline static void assign_double(double d, vector_type & v, const vect_dist_key_dx & key, const int (& comp)[3]) + { + pos_or_propL<vector_type,prop>::value(v,key)[comp[0]][comp[1]][comp[2]] = d; + } +}; + ///////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// template<unsigned int prp> diff --git a/src/Operators/Vector/vector_dist_operators.hpp b/src/Operators/Vector/vector_dist_operators.hpp index 1577ee2..2350fc6 100644 --- a/src/Operators/Vector/vector_dist_operators.hpp +++ b/src/Operators/Vector/vector_dist_operators.hpp @@ -1815,7 +1815,7 @@ public: o1.template value_nz<Sys_eqs>(p_map,key,cols,coeff,comp_ + var_id + comp[0]); } - inline vector_dist_expression_op<exp1,boost::mpl::int_<2>,VECT_COMP> operator[](int comp_) + inline vector_dist_expression_op<exp1,boost::mpl::int_<n+1>,VECT_COMP> operator[](int comp_) { int comp_n[n+1]; @@ -1823,7 +1823,7 @@ public: {comp_n[i] = comp[i];} comp_n[n] = comp_; - vector_dist_expression_op<exp1,boost::mpl::int_<2>,VECT_COMP> v_exp(o1,comp_n,var_id); + vector_dist_expression_op<exp1,boost::mpl::int_<n+1>,VECT_COMP> v_exp(o1,comp_n,var_id); return v_exp; } -- GitLab