Vlasiator ebf0dd394 on dev (v5.4.0 + 1054 commits)
Loading...
Searching...
No Matches
cpu_pitch_angle_diffusion.cpp
Go to the documentation of this file.
1/*
2 * This file is part of Vlasiator.
3 * Copyright 2010-2025 Finnish Meteorological Institute and University of Helsinki
4 *
5 * For details of usage, see the COPYING file and read the "Rules of the Road"
6 * at http://www.physics.helsinki.fi/vlasiator/
7 *
8 * This program is free software; you can redistribute it and/or modify
9 * it under the terms of the GNU General Public License as published by
10 * the Free Software Foundation; either version 2 of the License, or
11 * (at your option) any later version.
12 *
13 * This program is distributed in the hope that it will be useful,
14 * but WITHOUT ANY WARRANTY; without even the implied warranty of
15 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
16 * GNU General Public License for more details.
17 *
18 * You should have received a copy of the GNU General Public License along
19 * with this program; if not, write to the Free Software Foundation, Inc.,
20 * 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
21 */
22
23#include "../parameters.h"
24#include "../object_wrapper.h"
25#include <math.h>
26//#include <cmath> // NaN Inf checks
27#include <iostream>
28#include <fstream>
29#include <iomanip>
30#include <iterator>
31#include <Eigen/Geometry>
32#include "vec.h"
34
35#define MUSPACE(var,v_ind,mu_ind) var.at((mu_ind)*nbins_v + (v_ind))
36
37using namespace spatial_cell;
38using namespace Eigen;
39
40template <typename Lambda> inline static void loop_over_block(Lambda loop_body) {
41
42 for (int k = 0; k < WID; ++k) {
43 for (int j = 0; j < WID; j +=VECL/WID) { // Iterate through coordinates (z,y)
44
45 // create vectors with the i and j indices in the vector position on the plane.
46#if VECL == 4 && WID == 4
47 const Veci i_indices = Veci({0, 1, 2, 3});
48 const Veci j_indices = Veci({j, j, j, j});
49#elif VECL == 4 && WID == 8
50#error "__FILE__ : __LINE__ : VECL == 4 && WID == 8 cannot work!"
51#elif VECL == 8 && WID == 4
52 const Veci i_indices = Veci({0, 1, 2, 3,
53 0, 1, 2, 3});
54 const Veci j_indices = Veci({j, j, j, j,
55 j + 1, j + 1, j + 1, j + 1});
56#elif VECL == 8 && WID == 8
57 const Veci i_indices = Veci({0, 1, 2, 3, 4, 5, 6, 7});
58 const Veci j_indices = Veci({j, j, j, j, j, j, j, j});
59#elif VECL == 16 && WID == 4
60 const Veci i_indices = Veci({0, 1, 2, 3,
61 0, 1, 2, 3,
62 0, 1, 2, 3,
63 0, 1, 2, 3});
64 const Veci j_indices = Veci({j, j, j, j,
65 j + 1, j + 1, j + 1, j + 1,
66 j + 2, j + 2, j + 2, j + 2,
67 j + 3, j + 3, j + 3, j + 3});
68#elif VECL == 16 && WID == 8
69 const Veci i_indices = Veci({0, 1, 2, 3, 4, 5, 6, 7,
70 0, 1, 2, 3, 4, 5, 6, 7});
71 const Veci j_indices = Veci({j, j, j, j, j, j, j, j,
72 j+1, j+1, j+1, j+1, j+1, j+1, j+1, j+1});
73#elif VECL == 16 && WID == 16
74 const Veci i_indices = Veci({0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15});
75 const Veci j_indices = Veci({j, j, j, j, j, j, j, j, j, j, j, j, j, j, j, j});
76#elif VECL == 32 && WID == 4
77#error "__FILE__ : __LINE__ : VECL == 32 && WID == 4 cannot work, too long vector for one plane!"
78#elif VECL == 32 && WID == 8
79 const Veci i_indices = Veci({0, 1, 2, 3, 4, 5, 6, 7,
80 0, 1, 2, 3, 4, 5, 6, 7,
81 0, 1, 2, 3, 4, 5, 6, 7,
82 0, 1, 2, 3, 4, 5, 6, 7});
83 const Veci j_indices = Veci({j, j, j, j, j, j, j, j,
84 j+1, j+1, j+1, j+1, j+1, j+1, j+1, j+1,
85 j+2, j+2, j+2, j+2, j+2, j+2, j+2, j+2,
86 j+3, j+3, j+3, j+3, j+3, j+3, j+3, j+3});
87#elif VECL == 64 && WID == 4
88#error "__FILE__ : __LINE__ : VECL == 64 && WID == 4 cannot work, too long vector for one plane!"
89#elif VECL == 64 && WID == 8
90 const Veci i_indices = Veci({0, 1, 2, 3, 4, 5, 6, 7,
91 0, 1, 2, 3, 4, 5, 6, 7,
92 0, 1, 2, 3, 4, 5, 6, 7,
93 0, 1, 2, 3, 4, 5, 6, 7,
94 0, 1, 2, 3, 4, 5, 6, 7,
95 0, 1, 2, 3, 4, 5, 6, 7,
96 0, 1, 2, 3, 4, 5, 6, 7,
97 0, 1, 2, 3, 4, 5, 6, 7});
98 const Veci j_indices = Veci({j, j, j, j, j, j, j, j,
99 j+1, j+1, j+1, j+1, j+1, j+1, j+1, j+1,
100 j+2, j+2, j+2, j+2, j+2, j+2, j+2, j+2,
101 j+3, j+3, j+3, j+3, j+3, j+3, j+3, j+3,
102 j+4, j+4, j+4, j+4, j+4, j+4, j+4, j+4,
103 j+5, j+5, j+5, j+5, j+5, j+5, j+5, j+5,
104 j+6, j+6, j+6, j+6, j+6, j+6, j+6, j+6,
105 j+7, j+7, j+7, j+7, j+7, j+7, j+7, j+7});
106#else
107#error "This VECL && This WID cannot work!"
108#define xstr(s) str(s)
109#define str(s) #s
110#pragma message "VECL =" xstr(VECL)
111#pragma message "WID = "xstr(WID)
112#endif
113
114 loop_body(i_indices,j_indices,k);
115
116 }
117 }
118}
119
121 dccrg::Dccrg<SpatialCell,dccrg::Cartesian_Geometry>& mpiGrid,const uint popID){
122
123 // Ensure nu0 dat file is read, if requested
124 if (P::PADcoefficient < 0) {
126 }
127
130 const Real dmubins = 2.0/nbins_mu;
131
132 // resonance gap filling coefficient, not needed assuming even number of bins in mu-space
133 const Real epsilon = 0.0;
134
135 phiprof::Timer diffusionTimer {"pitch-angle-diffusion"};
136
137 const auto LocalCells=getLocalCells();
138 #pragma omp parallel
139 {
140 std::vector<int> fcount (nbins_v*nbins_mu,0); // Array to count number of f stored
141 std::vector<Realf> fmu (nbins_v*nbins_mu,0); // Array to store f(v,mu)
142 std::vector<Realf> dfdmu (nbins_v*nbins_mu,0); // Array to store dfdmu
143 std::vector<Realf> dfdmu2 (nbins_v*nbins_mu,0); // Array to store dfdmumu
144 std::vector<Realf> dfdt_mu(nbins_v*nbins_mu,0); // Array to store dfdt_mu
145 #pragma omp for
146 for (size_t CellIdx = 0; CellIdx < LocalCells.size(); CellIdx++) { // Iterate over all spatial cells
147
148 const auto CellID = LocalCells[CellIdx];
149 SpatialCell& cell = *mpiGrid[CellID];
150 const Real* parameters = cell.get_block_parameters(popID);
151 const vmesh::LocalID* nBlocks = cell.get_velocity_grid_length(popID);
152 const size_t meshID = getObjectWrapper().particleSpecies[popID].velocityMesh;
153 const vmesh::MeshParameters& vMesh = vmesh::getMeshWrapper()->velocityMeshes->at(meshID);
154
155 // Ensure mass conservation
156 Realf density_pre_adjust = 0.0;
157 Realf density_post_adjust = 0.0;
158 if (getObjectWrapper().particleSpecies[popID].sparse_conserve_mass) {
159 Vec vectorSum {0};
160 Vec vectorAdd {0};
161 for (size_t i=0; i<cell.get_number_of_velocity_blocks(popID)*WID3/VECL; ++i) {
162 vectorAdd.load(&cell.get_data(popID)[i*VECL]);
163 vectorSum += vectorAdd;
164 //density_pre_adjust += cell.get_data(popID)[i];
165 }
166 density_pre_adjust = horizontal_add(vectorSum);
167 }
168
169 Real dtTotalDiff = 0.0; // Diffusion time elapsed
170
171 const Real Vmax = 2*sqrt(3)*vMesh.meshLimits[1];
172 const Real dVbins = Vmax/nbins_v;
173
174 const Real bulkVX = cell.parameters[CellParams::VX];
175 const Real bulkVY = cell.parameters[CellParams::VY];
176 const Real bulkVZ = cell.parameters[CellParams::VZ];
177
178 bool currentSpatialLoopComplete;
179 Realf Sparsity;
180 std::array<Real,3> b;
181 Real nu0;
182
183 // Compute parameters
185 cell,
186 popID, CellIdx, currentSpatialLoopComplete,
187 Sparsity, b, nu0
188 );
189
190 // Enable nu0 disk output; skip cells where diffusion is not required (or diffusion coefficient is very small).
191 cell.parameters[CellParams::NU0] = nu0;
192 if (nu0 <= 0.001) {
193 continue;
194 }
195
196 while (dtTotalDiff < Parameters::dt) { // Substep loop
197
198 const Real RemainT = Parameters::dt - dtTotalDiff; //Remaining time before reaching simulation time step
199 Real checkCFL = std::numeric_limits<Real>::max();
200
201 // Initialised at each substep
202 std::fill(fmu.begin(), fmu.end(), 0.0);
203 std::fill(fcount.begin(), fcount.end(), 0);
204
205 // Build 2d array of f(v,mu)
206 for (vmesh::LocalID n=0; n<cell.get_number_of_velocity_blocks(popID); n++) { // Iterate through velocity blocks
207
208 loop_over_block([&](Veci i_indices, Veci j_indices, int k) -> void { // Lambda function processor
209
210 //Get velocity space coordinates
212 + (to_realf(i_indices) + 0.5)*parameters[n * BlockParams::N_VELOCITY_BLOCK_PARAMS + BlockParams::DVX]);
214 + (to_realf(j_indices) + 0.5)*parameters[n * BlockParams::N_VELOCITY_BLOCK_PARAMS + BlockParams::DVY]);
216 + (k + 0.5)*parameters[n * BlockParams::N_VELOCITY_BLOCK_PARAMS + BlockParams::DVZ]);
217
218 const Vec VplasmaX = VX - bulkVX;
219 const Vec VplasmaY = VY - bulkVY;
220 const Vec VplasmaZ = VZ - bulkVZ;
221
223 const Vec Vpara = VplasmaX*b[0] + VplasmaY*b[1] + VplasmaZ*b[2];
224 const Vec mu = Vpara/(normV+std::numeric_limits<Real>::min()); // + min value to avoid division by 0.
225
226 const Veci Vindex = roundi(floor((normV) / dVbins));
227 const Vec Vmu = dVbins * (to_realf(Vindex)+0.5); // Take value at the center of the mu cell
228 Veci muindex = roundi(floor((mu+1.0) / dmubins));
229
230 Vec CellValue;
231 CellValue.load(&cell.get_data(n,popID)[WID2*k + WID*j_indices[0] + i_indices[0]]);
232 const Vec increment = 2.0 * M_PI * Vmu*Vmu * CellValue;
233 for (uint i = 0; i<VECL; i++) {
234 // Safety check to handle edge case where mu = exactly 1.0
235 const int mui = std::max(0,std::min((int)muindex[i],nbins_mu-1));
236 const int vi = std::max(0,std::min((int)Vindex[i],nbins_v-1));
237 MUSPACE(fmu,vi,mui) += increment[i];
238 MUSPACE(fcount,vi,mui) += 1;
239 }
240 }); // End of Lambda
241 } // End blocks
242
243 int cRight;
244 int cLeft;
245
246 // Compute space/time derivatives (take first non-zero neighbours) & CFL & Ddt
247 for (int indv = 0; indv < nbins_v; indv++) {
248 const Real Vmu = dVbins * (float(indv)+0.5);
249
250 // Divide f by count (independent of v but needs to be computed for all mu before derivatives)
251 for(int indmu = 0; indmu < nbins_mu; indmu++) {
252 if (MUSPACE(fcount,indv,indmu) == 0 || MUSPACE(fmu,indv,indmu) <= 0.0) {
253 MUSPACE(fmu,indv,indmu) = 0;
254 } else {
255 MUSPACE(fmu,indv,indmu) = MUSPACE(fmu,indv,indmu) / MUSPACE(fcount,indv,indmu);
256 }
257 }
258
259 // Search limits for how many cells in mu-direction should be max evaluated when searching for a near neighbour?
260 // Assuming some oversampling; changing these values may result in method breaking at very small plasma frame velocities.
261 const int rlimit = nbins_mu-1;
262 const int llimit = 0;
263
264 for(int indmu = 0; indmu < nbins_mu; indmu++) {
265 // Compute spatial derivatives
266 if (indmu == 0) {
267 cLeft = 0;
268 cRight = 1;
269 while( (MUSPACE(fcount,indv,indmu + cRight) == 0) && (indmu + cRight < rlimit) ) { cRight += 1; }
270 if( (MUSPACE(fcount,indv,indmu + cRight) == 0) && (indmu + cRight == rlimit) ) { cRight = 0; }
271 } else if (indmu == nbins_mu-1) {
272 cLeft = 1;
273 cRight = 0;
274 while( (MUSPACE(fcount,indv,indmu - cLeft) == 0) && (indmu - cLeft > llimit) ) { cLeft += 1; }
275 if( (MUSPACE(fcount,indv,indmu - cLeft) == 0) && (indmu - cLeft == llimit) ) { cLeft = 0; }
276 } else {
277 cLeft = 1;
278 cRight = 1;
279 while( (MUSPACE(fcount,indv,indmu + cRight) == 0) && (indmu + cRight < rlimit) ) { cRight += 1; }
280 if( (MUSPACE(fcount,indv,indmu + cRight) == 0) && (indmu + cRight == rlimit) ) { cRight = 0; }
281 while( (MUSPACE(fcount,indv,indmu - cLeft ) == 0) && (indmu - cLeft > llimit) ) { cLeft += 1; }
282 if( (MUSPACE(fcount,indv,indmu - cLeft ) == 0) && (indmu - cLeft == llimit) ) { cLeft = 0; }
283 }
284 if( (cRight == 0) && (cLeft != 0) ) {
285 MUSPACE(dfdmu ,indv,indmu) = (MUSPACE(fmu,indv,indmu + cRight) - MUSPACE(fmu,indv,indmu - cLeft))/((cRight + cLeft)*dmubins) ;
286 MUSPACE(dfdmu2,indv,indmu) = 0.0;
287 } else if( (cLeft == 0) && (cRight != 0) ) {
288 MUSPACE(dfdmu ,indv,indmu) = (MUSPACE(fmu,indv,indmu + cRight) - MUSPACE(fmu,indv,indmu - cLeft))/((cRight + cLeft)*dmubins) ;
289 MUSPACE(dfdmu2,indv,indmu) = 0.0;
290 } else if( (cLeft == 0) && (cRight == 0) ) {
291 MUSPACE(dfdmu ,indv,indmu) = 0.0;
292 MUSPACE(dfdmu2,indv,indmu) = 0.0;
293 } else {
294 MUSPACE(dfdmu ,indv,indmu) = ( MUSPACE(fmu,indv,indmu + cRight) - MUSPACE(fmu,indv,indmu - cLeft))/((cRight + cLeft)*dmubins) ;
295 MUSPACE(dfdmu2,indv,indmu) = ( (MUSPACE(fmu,indv,indmu + cRight) - MUSPACE(fmu,indv,indmu))/(cRight*dmubins) - (MUSPACE(fmu,indv,indmu) - MUSPACE(fmu,indv,indmu - cLeft))/(cLeft*dmubins) ) / (0.5 * dmubins * (cRight + cLeft));
296 }
297
298 // Compute time derivative
299 const Realf mu = (indmu+0.5)*dmubins - 1.0;
300 const Realf Dmumu = nu0/2.0 * ( abs(mu)/(1.0 + abs(mu)) + epsilon ) * (1.0 - mu*mu);
301 const Realf dDmu = nu0/2.0 * ( (mu/abs(mu)) * ((1.0 - mu*mu)/((1.0 + abs(mu))*(1.0 + abs(mu)))) - 2.0*mu*( abs(mu)/(1.0 + abs(mu)) + epsilon));
302 // We divide dfdt_mu by the normalization factor 2pi*v^2 already here.
303 const Realf dfdt_mu_val = ( dDmu * MUSPACE(dfdmu,indv,indmu) + Dmumu * MUSPACE(dfdmu2,indv,indmu) ) / (2.0 * M_PI * Vmu*Vmu);
304 MUSPACE(dfdt_mu,indv,indmu) = dfdt_mu_val;
305
306 // Only consider CFL for non-negative phase-space cells above the sparsity threshold
307 const Realf CellValue = MUSPACE(fmu,indv,indmu) / (2.0 * M_PI * Vmu*Vmu);
308 const Realf absdfdt = abs(MUSPACE(dfdt_mu,indv,indmu)); // Already scaled
309 if (absdfdt > 0.0 && CellValue > Sparsity) {
310 checkCFL = std::min(CellValue * Parameters::PADCFL * (1.0/absdfdt), checkCFL);
311 }
312 } // End mu loop
313 } // End v loop
314
315 // Compute Ddt
316 Real Ddt = checkCFL;
317 if (Ddt > RemainT) {
318 Ddt = RemainT;
319 }
320 dtTotalDiff = dtTotalDiff + Ddt;
321
322 for (vmesh::LocalID n=0; n<cell.get_number_of_velocity_blocks(popID); n++) { // Iterate through velocity blocks
323
324 loop_over_block([&](Veci i_indices, Veci j_indices, int k) -> void { // Lambda function processor
325
326 //Get velocity space coordinates
328 + (to_realf(i_indices) + 0.5)*parameters[n * BlockParams::N_VELOCITY_BLOCK_PARAMS + BlockParams::DVX]);
330 + (to_realf(j_indices) + 0.5)*parameters[n * BlockParams::N_VELOCITY_BLOCK_PARAMS + BlockParams::DVY]);
332 + (k + 0.5)*parameters[n * BlockParams::N_VELOCITY_BLOCK_PARAMS + BlockParams::DVZ]);
333
334 const Vec VplasmaX = VX - bulkVX;
335 const Vec VplasmaY = VY - bulkVY;
336 const Vec VplasmaZ = VZ - bulkVZ;
337
339 const Vec Vpara = VplasmaX*b[0] + VplasmaY*b[1] + VplasmaZ*b[2];
340 const Vec mu = Vpara/(normV+std::numeric_limits<Real>::min()); // + min value to avoid division by 0.
341
342 const Veci Vindex = roundi(floor((normV) / dVbins));
343 const Vec Vmu = dVbins * (to_realf(Vindex)+0.5); // Take value at the center of the mu cell
344 Veci muindex = roundi(floor((mu+1.0) / dmubins));
345
346 // Compute dfdt
347 std::array<Realf,VECL> dfdt = {0};
348 for (uint i = 0; i < VECL; i++) {
349 // Safety check to handle edge case where mu = exactly 1.0
350 const int mui = std::max(0,std::min((int)muindex[i],nbins_mu-1));
351 const int vi = std::max(0,std::min((int)Vindex[i],nbins_v-1));
352 dfdt[i] = MUSPACE(dfdt_mu,vi,mui); // dfdt_mu was scaled back down by 2pi*v^2 on creation
353 }
354 Vec dfdtUpdate;
355 dfdtUpdate.load(&dfdt[0]);
356
357 // Update cell value, ensuring result is non-negative
358 Vec CellValue;
359 CellValue.load(&cell.get_data(n,popID)[WID2*k + WID*j_indices[0] + i_indices[0]]);
360 Vec NewCellValue = CellValue + dfdtUpdate * Ddt;
361 const Vecb lessZero = NewCellValue < 0.0;
363 NewCellValue.store(&cell.get_data(n,popID)[WID2*k + WID*j_indices[0] + i_indices[0]]);
364 }); // End of Lambda
365 } // End Blocks
366
367 } // End Time loop
368
369 // Ensure mass conservation
370 if (getObjectWrapper().particleSpecies[popID].sparse_conserve_mass) {
371 Vec vectorSum {0};
372 Vec vectorAdd {0};
373 for (size_t i=0; i<cell.get_number_of_velocity_blocks(popID)*WID3/VECL; ++i) {
374 vectorAdd.load(&cell.get_data(popID)[i*VECL]);
375 vectorSum += vectorAdd;
376 }
377 density_post_adjust = horizontal_add(vectorSum);
378
379 if (density_post_adjust != 0.0 && density_pre_adjust != density_post_adjust) {
380 const Vec adjustRatio = density_pre_adjust/density_post_adjust;
381 Vec vectorAdjust;
382 for (size_t i=0; i<cell.get_number_of_velocity_blocks(popID)*WID3/VECL; ++i) {
383 vectorAdjust.load(&cell.get_data(popID)[i*VECL]);
384 vectorAdjust *= adjustRatio;
385 vectorAdjust.store(&cell.get_data(popID)[i*VECL]);
386 }
387 }
388 }
389 } // End spatial cell loop
390 } // End parallel workshare region
391} // End function
for i
Definition Dispersion.m:24
sqrt(1.0+vA *vA/(c *c))) % Ion-acoustic wave cS
vmesh::LocalID get_number_of_velocity_blocks(const uint popID) const
Real * get_block_parameters(const uint popID)
const vmesh::LocalID * get_velocity_grid_length(const uint popID)
Realf * get_data(const uint popID)
std::array< Real, CellParams::N_SPATIAL_CELL_PARAMS > parameters
const std::vector< CellID > & getLocalCells()
Definition main.cpp:39
#define WID
Definition common.h:514
const int WID3
Definition common.h:517
const int WID2
Definition common.h:516
void readNuArrayFromFile()
void computePitchAngleDiffusionParameters(SpatialCell &cell, const uint popID, const size_t CellIdx, bool &currentSpatialLoopComplete, Realf &sparsity, std::array< Real, 3 > &b, Real &nu0)
static void loop_over_block(Lambda loop_body)
void pitchAngleDiffusion(dccrg::Dccrg< SpatialCell, dccrg::Cartesian_Geometry > &mpiGrid, const uint popID)
#define MUSPACE(var, v_ind, mu_ind)
float Real
Definition definitions.h:41
uint64_t CellID
Definition definitions.h:54
float Realf
Definition definitions.h:33
const Real VY
__global__ void size_t const vmesh::VelocityBlockContainer *__restrict__ const Real Real Real Real Realf int Real const Real int int nbins_mu
const Real increment
const int Vindex
__global__ void size_t const vmesh::VelocityBlockContainer *__restrict__ const Real Real Real Real Realf int Real dVbins
const Real VplasmaZ
const Real mu
const int vi
const Real VplasmaY
__global__ void size_t const vmesh::VelocityBlockContainer *__restrict__ const Real Real Real Real Realf int Real const Real int nbins_v
__global__ void size_t const vmesh::VelocityBlockContainer *__restrict__ const Real Real Real Real Realf int Real const Real dmubins
const int mui
const Real VplasmaX
const Real VZ
const int j
const bool lessZero
const Real Vmu
const Real Vpara
const Real VX
const int k
const Real normV
ObjectWrapper & getObjectWrapper()
Definition main.cpp:33
@ N_VELOCITY_BLOCK_PARAMS
Definition common.h:115
uint32_t LocalID
Definition definitions.h:60
ARCH_HOSTDEV MeshWrapper * getMeshWrapper()
std::vector< species::Species > particleSpecies
static int PADvbins
Definition parameters.h:240
static Realf PADcoefficient
Definition parameters.h:238
static Realf PADCFL
Definition parameters.h:239
static int PADmubins
Definition parameters.h:241
static Real dt
Definition parameters.h:55
std::array< vmesh::MeshParameters, MAX_VMESH_PARAMETERS_COUNT > * velocityMeshes
An interface to a type with floating point values.
static ARCH_HOSTDEV VecSimple< int > roundi(VecSimple< T > const &a)
static ARCH_HOSTDEV VecSimple< T > select(VecSimple< bool > const &a, VecSimple< T > const &b, VecSimple< T > const &c)
static ARCH_HOSTDEV T horizontal_add(VecSimple< T > const &a)
static ARCH_HOSTDEV VecSimple< T > abs(const VecSimple< T > &l)
static ARCH_HOSTDEV VecSimple< T > floor(VecSimple< T > const &a)