Vlasiator ebf0dd394 on dev (v5.4.0 + 1054 commits)
Loading...
Searching...
No Matches
gpu_pitch_angle_diffusion.cpp File Reference
#include "../parameters.h"
#include "../object_wrapper.h"
#include <math.h>
#include <iostream>
#include <fstream>
#include <iomanip>
#include <iterator>
#include <Eigen/Geometry>
#include "common_pitch_angle_diffusion.hpp"
#include "../arch/gpu_base.hpp"
#include "../spatial_cells/block_adjust_gpu.hpp"
Include dependency graph for gpu_pitch_angle_diffusion.cpp:

Go to the source code of this file.

Macros

#define GPUCELLMUSPACE(var, cellIdx, v_ind, mu_ind)

Functions

__global__ void __launch_bounds__ (WID3) build2dArrayOfFvmu_kernel(size_t *dev_cellIdxArray
atomicAdd & GPUCELLMUSPACE (dev_fmu, cellIdx, vi, mui), increment
atomicAdd & GPUCELLMUSPACE (dev_fcount, cellIdx, vi, mui), 1
dev_velocityBlockContainer[cellIdx] getData ()[velocityIdx *WID3+k *WID2+j *WID+i]
__global__ void computeDerivativesCFLDdt_kernel (size_t *dev_smallCellIdxArray, int *dev_fcount, Realf *dev_fmu, Real *dev_nu0Values, Realf *dev_dfdt_mu, int *dev_cellIdxKeys, Real *dev_potentialDdtValues, Realf *dev_sparsity, Real dVbins, const Real dmubins, const Real epsilon, Realf PADCFL, int nbins_v, int nbins_mu, int blocksPerSpatialCell, int lastBlockSize)
__global__ void reduceDdtValues_kernel (int *dev_cellIdxKeys, Real *dev_potentialDdtValues, Real *dev_Ddt, int blocksPerSpatialCell)
__global__ void __launch_bounds__ (Hashinator::defaults::MAX_BLOCKSIZE/2) dividefByCount_kernel(size_t *dev_smallCellIdxArray
 if (idx >=maxThreadIndex)
 if (GPUCELLMUSPACE(dev_fcount, cellIdx, indv, indmu)==0||GPUCELLMUSPACE(dev_fmu, cellIdx, indv, indmu)<=0.0)
 if (totalBlockIndex >=(size_t) numberOfComputedVelocityBlocks)
 while (left<=right)
 for (uint velocityIdx=0;velocityIdx< numberOfVelocityCells;velocityIdx++)
 __syncthreads ()
 for (int s=blockSize/2;s > 0;s > >=1)
 if (threadIndex==0)
 if (dev_densityPostAdjust[cellIdx]==0.0||dev_densityPreAdjust[cellIdx]==dev_densityPostAdjust[cellIdx])
void pitchAngleDiffusion (dccrg::Dccrg< spatial_cell::SpatialCell, dccrg::Cartesian_Geometry > &mpiGrid, const uint popID)

Variables

__global__ void size_t * dev_velocityIdxArray = totalBlockIndex - dev_cellIdxStartCutoff[cellIndex]
__global__ void size_t const vmesh::VelocityBlockContainer *__restrict__ const * dev_velocityBlockContainer
__global__ void size_t const vmesh::VelocityBlockContainer *__restrict__ const Real * dev_bulkVX
__global__ void size_t const vmesh::VelocityBlockContainer *__restrict__ const Real Real * dev_bulkVY
__global__ void size_t const vmesh::VelocityBlockContainer *__restrict__ const Real Real Real * dev_bulkVZ
__global__ void size_t const vmesh::VelocityBlockContainer *__restrict__ const Real Real Real Real * dev_bValues
__global__ void size_t const vmesh::VelocityBlockContainer *__restrict__ const Real Real Real Real Realf * dev_fmu
__global__ void size_t const vmesh::VelocityBlockContainer *__restrict__ const Real Real Real Real Realf int * dev_fcount
__global__ void size_t const vmesh::VelocityBlockContainer *__restrict__ const Real Real Real Real Realf int Real dVbins
__global__ void size_t const vmesh::VelocityBlockContainer *__restrict__ const Real Real Real Real Realf int Real const Real dmubins
__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 int int nbins_mu
const int i = threadIdx.x
const int j = threadIdx.y
const int k = threadIdx.z
size_t cellIdx = dev_cellIdxArray[totalBlockIndex]
size_t velocityIdx = dev_velocityIdxArray[totalBlockIndex]
const Real *__restrict__ blockParameters = dev_velocityBlockContainer[cellIdx]->getParameters(velocityIdx)
const Real VX = blockParameters[BlockParams::VXCRD] + (i + 0.5)*blockParameters[BlockParams::DVX]
const Real VY = blockParameters[BlockParams::VYCRD] + (j + 0.5)*blockParameters[BlockParams::DVY]
const Real VZ = blockParameters[BlockParams::VZCRD] + (k + 0.5)*blockParameters[BlockParams::DVZ]
const Real VplasmaX = VX - dev_bulkVX[cellIdx]
const Real VplasmaY = VY - dev_bulkVY[cellIdx]
const Real VplasmaZ = VZ - dev_bulkVZ[cellIdx]
const Real normV = sqrt(VplasmaX*VplasmaX + VplasmaY*VplasmaY + VplasmaZ*VplasmaZ)
const Real Vpara = VplasmaX*dev_bValues[3*cellIdx] + VplasmaY*dev_bValues[3*cellIdx+1] + VplasmaZ*dev_bValues[3*cellIdx+2]
const Real mu = Vpara/(normV+std::numeric_limits<Real>::min())
const int Vindex = static_cast<int>(std::nearbyint(floor((normV) / dVbins)))
const Real Vmu = dVbins * (Vindex+0.5)
int muindex = static_cast<int>(std::nearbyint(floor((mu+1.0) / dmubins)))
const Realf cellValue = dev_velocityBlockContainer[cellIdx]->getData()[velocityIdx*WID3+k*WID2+j*WID+i]
const Real increment = 2.0 * M_PI * Vmu*Vmu * cellValue
const int mui = std::max(0,std::min(muindex,nbins_mu-1))
const int vi = std::max(0,std::min(Vindex,nbins_v-1))
__global__ void size_t * dev_remappedCellIdxArray
__global__ void size_t size_t vmesh::VelocityBlockContainer *__restrict__ Real Real Real Real Realf * dev_dfdt_mu
__global__ void size_t size_t vmesh::VelocityBlockContainer *__restrict__ Real Real Real Real Realf Real * dev_Ddt
Realf dfdt = 0.0
Realf NewCellValue = (dev_velocityBlockContainer[cellIdx]->getData()[velocityIdx*WID3+k*WID2+j*WID+i]) + dfdt * dev_Ddt[cellIdx]
const bool lessZero = (NewCellValue < 0.0)
__global__ void Realf int int int int maxThreadIndex
int indv = idx%nbins_v
int indmu = (idx/nbins_v)%nbins_mu
int spatialBlockIndex = idx/(nbins_v*nbins_mu)
 else
__global__ void size_t size_t * dev_cellIdxStartCutoff
__global__ void size_t size_t int numberOfComputedVelocityBlocks
__global__ void size_t size_t int int maxCellIndex
int left = 0
int right = maxCellIndex - 1
int cellIndex = 0
 dev_cellIdxArray [totalBlockIndex] = cellIndex
const int threadIndex = k*WID2+j*WID+i
const int blockSize = blockDim.x*blockDim.y*blockDim.z
__shared__ Realf localDensity [WID3] = 0.0
const uint numberOfVelocityCells = dev_velocityBlockContainer[cellIdx]->size()
__global__ void Realf * dev_densityPostAdjust
Realf adjustRatio = dev_densityPreAdjust[cellIdx]/dev_densityPostAdjust[cellIdx]

Macro Definition Documentation

◆ GPUCELLMUSPACE

#define GPUCELLMUSPACE ( var,
cellIdx,
v_ind,
mu_ind )
Value:
var[(cellIdx)*nbins_v*nbins_mu+(mu_ind)*nbins_v + (v_ind)]
nbins_mu
__global__ void size_t const vmesh::VelocityBlockContainer *__restrict__ const Real Real Real Real Realf int Real const Real int int nbins_mu
Definition gpu_pitch_angle_diffusion.cpp:55
nbins_v
__global__ void size_t const vmesh::VelocityBlockContainer *__restrict__ const Real Real Real Real Realf int Real const Real int nbins_v
Definition gpu_pitch_angle_diffusion.cpp:53
cellIdx
size_t cellIdx
Definition gpu_pitch_angle_diffusion.cpp:62

Definition at line 39 of file gpu_pitch_angle_diffusion.cpp.

Function Documentation

◆ __launch_bounds__() [1/2]

__global__ void __launch_bounds__ ( Hashinator::defaults::MAX_BLOCKSIZE/ 2)

◆ __launch_bounds__() [2/2]

__global__ void __launch_bounds__ ( WID3 )

◆ __syncthreads()

__syncthreads ( )
Here is the caller graph for this function:

◆ computeDerivativesCFLDdt_kernel()

__global__ void computeDerivativesCFLDdt_kernel ( size_t * dev_smallCellIdxArray,
int * dev_fcount,
Realf * dev_fmu,
Real * dev_nu0Values,
Realf * dev_dfdt_mu,
int * dev_cellIdxKeys,
Real * dev_potentialDdtValues,
Realf * dev_sparsity,
Real dVbins,
const Real dmubins,
const Real epsilon,
Realf PADCFL,
int nbins_v,
int nbins_mu,
int blocksPerSpatialCell,
int lastBlockSize )

Definition at line 151 of file gpu_pitch_angle_diffusion.cpp.

Here is the call graph for this function:
Here is the caller graph for this function:

◆ for() [1/2]

for ( int s = blockSize / 2; s,
0;s > > = 1 )

Definition at line 405 of file gpu_pitch_angle_diffusion.cpp.

Here is the call graph for this function:

◆ for() [2/2]

for ( )

Definition at line 399 of file gpu_pitch_angle_diffusion.cpp.

◆ getData()

dev_velocityBlockContainer[cellIdx] getData ( )
Here is the caller graph for this function:

◆ GPUCELLMUSPACE() [1/2]

atomicAdd & GPUCELLMUSPACE ( dev_fcount ,
cellIdx ,
vi ,
mui  )

◆ GPUCELLMUSPACE() [2/2]

atomicAdd & GPUCELLMUSPACE ( dev_fmu ,
cellIdx ,
vi ,
mui  )

◆ if() [1/5]

if ( dev_densityPostAdjust [cellIdx] = = 0.0 || dev_densityPreAdjust[cellIdx] == dev_densityPostAdjust[cellIdx])

Definition at line 429 of file gpu_pitch_angle_diffusion.cpp.

◆ if() [2/5]

if ( GPUCELLMUSPACE(dev_fcount, cellIdx, indv, indmu) = = 0 || GPUCELLMUSPACE(dev_fmu,cellIdx,indv,indmu) <= 0.0)

Definition at line 340 of file gpu_pitch_angle_diffusion.cpp.

◆ if() [3/5]

if ( idx >= maxThreadIndex)

Definition at line 333 of file gpu_pitch_angle_diffusion.cpp.

◆ if() [4/5]

if ( threadIndex = = 0)

Definition at line 413 of file gpu_pitch_angle_diffusion.cpp.

◆ if() [5/5]

if ( totalBlockIndex >=(size_t) numberOfComputedVelocityBlocks)

Definition at line 357 of file gpu_pitch_angle_diffusion.cpp.

◆ pitchAngleDiffusion()

void pitchAngleDiffusion ( dccrg::Dccrg< spatial_cell::SpatialCell, dccrg::Cartesian_Geometry > & mpiGrid,
const uint popID )

Definition at line 440 of file gpu_pitch_angle_diffusion.cpp.

Here is the call graph for this function:

◆ reduceDdtValues_kernel()

__global__ void reduceDdtValues_kernel ( int * dev_cellIdxKeys,
Real * dev_potentialDdtValues,
Real * dev_Ddt,
int blocksPerSpatialCell )

Definition at line 279 of file gpu_pitch_angle_diffusion.cpp.

Here is the call graph for this function:
Here is the caller graph for this function:

◆ while()

while ( left<= right)

Definition at line 368 of file gpu_pitch_angle_diffusion.cpp.

Variable Documentation

◆ adjustRatio

Realf adjustRatio = dev_densityPreAdjust[cellIdx]/dev_densityPostAdjust[cellIdx]

Definition at line 431 of file gpu_pitch_angle_diffusion.cpp.

◆ blockParameters

const Real *__restrict__ blockParameters = dev_velocityBlockContainer[cellIdx]->getParameters(velocityIdx)

Definition at line 65 of file gpu_pitch_angle_diffusion.cpp.

◆ blockSize

const int blockSize = blockDim.x*blockDim.y*blockDim.z

Definition at line 392 of file gpu_pitch_angle_diffusion.cpp.

◆ cellIdx

const int cellIdx = dev_cellIdxArray[totalBlockIndex]

Definition at line 62 of file gpu_pitch_angle_diffusion.cpp.

◆ cellIndex

int cellIndex = 0

Definition at line 362 of file gpu_pitch_angle_diffusion.cpp.

◆ cellValue

const Realf cellValue = dev_velocityBlockContainer[cellIdx]->getData()[velocityIdx*WID3+k*WID2+j*WID+i]

Definition at line 84 of file gpu_pitch_angle_diffusion.cpp.

◆ dev_bulkVX

__global__ void size_t size_t vmesh::VelocityBlockContainer *__restrict__ Real * dev_bulkVX

Definition at line 45 of file gpu_pitch_angle_diffusion.cpp.

◆ dev_bulkVY

__global__ void size_t size_t vmesh::VelocityBlockContainer *__restrict__ Real Real * dev_bulkVY

Definition at line 46 of file gpu_pitch_angle_diffusion.cpp.

◆ dev_bulkVZ

__global__ void size_t size_t vmesh::VelocityBlockContainer *__restrict__ Real Real Real * dev_bulkVZ

Definition at line 47 of file gpu_pitch_angle_diffusion.cpp.

◆ dev_bValues

__global__ void size_t size_t vmesh::VelocityBlockContainer *__restrict__ Real Real Real Real * dev_bValues

Definition at line 48 of file gpu_pitch_angle_diffusion.cpp.

◆ dev_cellIdxArray

dev_cellIdxArray[totalBlockIndex] = cellIndex

Definition at line 378 of file gpu_pitch_angle_diffusion.cpp.

◆ dev_cellIdxStartCutoff

__global__ void size_t size_t* dev_cellIdxStartCutoff

Definition at line 350 of file gpu_pitch_angle_diffusion.cpp.

◆ dev_Ddt

__global__ void size_t size_t vmesh::VelocityBlockContainer* __restrict__ Real Real Real Real Realf Real* dev_Ddt

Definition at line 105 of file gpu_pitch_angle_diffusion.cpp.

◆ dev_densityPostAdjust

__global__ void Realf* dev_densityPostAdjust

Definition at line 420 of file gpu_pitch_angle_diffusion.cpp.

◆ dev_dfdt_mu

__global__ void size_t size_t vmesh::VelocityBlockContainer* __restrict__ Real Real Real Real Realf* dev_dfdt_mu

Definition at line 104 of file gpu_pitch_angle_diffusion.cpp.

◆ dev_fcount

__global__ void Realf int * dev_fcount

Definition at line 50 of file gpu_pitch_angle_diffusion.cpp.

◆ dev_fmu

__global__ void Realf * dev_fmu

Definition at line 49 of file gpu_pitch_angle_diffusion.cpp.

◆ dev_remappedCellIdxArray

__global__ void size_t* dev_remappedCellIdxArray

Definition at line 97 of file gpu_pitch_angle_diffusion.cpp.

◆ dev_velocityBlockContainer

__global__ void Realf vmesh::VelocityBlockContainer *__restrict__ * dev_velocityBlockContainer
Initial value:
{
const int i = threadIdx.x
i
for i
Definition Dispersion.m:24

Definition at line 44 of file gpu_pitch_angle_diffusion.cpp.

◆ dev_velocityIdxArray

__global__ void size_t * dev_velocityIdxArray = totalBlockIndex - dev_cellIdxStartCutoff[cellIndex]

Definition at line 43 of file gpu_pitch_angle_diffusion.cpp.

◆ dfdt

dfdt = 0.0

Definition at line 138 of file gpu_pitch_angle_diffusion.cpp.

◆ dmubins

__global__ void size_t size_t vmesh::VelocityBlockContainer *__restrict__ Real Real Real Real Realf Real Real const Real dmubins

Definition at line 52 of file gpu_pitch_angle_diffusion.cpp.

◆ dVbins

__global__ void size_t size_t vmesh::VelocityBlockContainer *__restrict__ Real Real Real Real Realf Real Real dVbins

Definition at line 51 of file gpu_pitch_angle_diffusion.cpp.

◆ else

else
Initial value:
{
GPUCELLMUSPACE(dev_fmu,cellIdx,indv,indmu) = GPUCELLMUSPACE(dev_fmu,cellIdx,indv,indmu) / GPUCELLMUSPACE(dev_fcount,cellIdx,indv,indmu)
GPUCELLMUSPACE
#define GPUCELLMUSPACE(var, cellIdx, v_ind, mu_ind)
Definition gpu_pitch_angle_diffusion.cpp:39
indmu
int indmu
Definition gpu_pitch_angle_diffusion.cpp:336
dev_fmu
__global__ void size_t const vmesh::VelocityBlockContainer *__restrict__ const Real Real Real Real Realf * dev_fmu
Definition gpu_pitch_angle_diffusion.cpp:49
dev_fcount
__global__ void size_t const vmesh::VelocityBlockContainer *__restrict__ const Real Real Real Real Realf int * dev_fcount
Definition gpu_pitch_angle_diffusion.cpp:50
indv
int indv
Definition gpu_pitch_angle_diffusion.cpp:335

Definition at line 342 of file gpu_pitch_angle_diffusion.cpp.

◆ i

const int i = threadIdx.x

Definition at line 59 of file gpu_pitch_angle_diffusion.cpp.

◆ increment

const Real increment = 2.0 * M_PI * Vmu*Vmu * cellValue

Definition at line 85 of file gpu_pitch_angle_diffusion.cpp.

◆ indmu

int indmu = (idx/nbins_v)%nbins_mu

Definition at line 336 of file gpu_pitch_angle_diffusion.cpp.

◆ indv

int indv = idx%nbins_v

Definition at line 335 of file gpu_pitch_angle_diffusion.cpp.

◆ j

const int j = threadIdx.y

Definition at line 60 of file gpu_pitch_angle_diffusion.cpp.

◆ k

const int k = threadIdx.z

Definition at line 61 of file gpu_pitch_angle_diffusion.cpp.

◆ left

int left = 0

Definition at line 360 of file gpu_pitch_angle_diffusion.cpp.

◆ lessZero

const bool lessZero = (NewCellValue < 0.0)

Definition at line 146 of file gpu_pitch_angle_diffusion.cpp.

◆ localDensity

localDensity[threadIndex] = 0.0

Definition at line 394 of file gpu_pitch_angle_diffusion.cpp.

◆ maxCellIndex

__global__ void size_t size_t int int maxCellIndex
Initial value:
{
size_t totalBlockIndex = blockIdx.x*blockDim.x + threadIdx.x

Definition at line 352 of file gpu_pitch_angle_diffusion.cpp.

◆ maxThreadIndex

__global__ void Realf int int int int maxThreadIndex
Initial value:
{
int idx = blockIdx.x*blockDim.x + threadIdx.x

Definition at line 328 of file gpu_pitch_angle_diffusion.cpp.

◆ mu

const Real mu = Vpara/(normV+std::numeric_limits<Real>::min())

Definition at line 78 of file gpu_pitch_angle_diffusion.cpp.

◆ mui

const int mui = std::max(0,std::min(muindex,nbins_mu-1))

Definition at line 87 of file gpu_pitch_angle_diffusion.cpp.

◆ muindex

int muindex = static_cast<int>(std::nearbyint(floor((mu+1.0) / dmubins)))

Definition at line 82 of file gpu_pitch_angle_diffusion.cpp.

◆ nbins_mu

__global__ void Realf int int int nbins_mu
Initial value:
{
int totalBlockIndex = blockIdx.x

Definition at line 54 of file gpu_pitch_angle_diffusion.cpp.

◆ nbins_v

__global__ void Realf int int nbins_v

Definition at line 53 of file gpu_pitch_angle_diffusion.cpp.

◆ NewCellValue

NewCellValue = (dev_velocityBlockContainer[cellIdx]->getData()[velocityIdx*WID3+k*WID2+j*WID+i]) + dfdt * dev_Ddt[cellIdx]

Definition at line 145 of file gpu_pitch_angle_diffusion.cpp.

◆ normV

const Real normV = sqrt(VplasmaX*VplasmaX + VplasmaY*VplasmaY + VplasmaZ*VplasmaZ)

Definition at line 76 of file gpu_pitch_angle_diffusion.cpp.

◆ numberOfComputedVelocityBlocks

__global__ void size_t size_t int numberOfComputedVelocityBlocks

Definition at line 351 of file gpu_pitch_angle_diffusion.cpp.

◆ numberOfVelocityCells

const uint numberOfVelocityCells = dev_velocityBlockContainer[cellIdx]->size()

Definition at line 397 of file gpu_pitch_angle_diffusion.cpp.

◆ right

int right = maxCellIndex - 1

Definition at line 361 of file gpu_pitch_angle_diffusion.cpp.

◆ spatialBlockIndex

int spatialBlockIndex = idx/(nbins_v*nbins_mu)

Definition at line 337 of file gpu_pitch_angle_diffusion.cpp.

◆ threadIndex

const int threadIndex = k*WID2+j*WID+i

Definition at line 390 of file gpu_pitch_angle_diffusion.cpp.

◆ velocityIdx

size_t velocityIdx = dev_velocityIdxArray[totalBlockIndex]

Definition at line 63 of file gpu_pitch_angle_diffusion.cpp.

◆ vi

const int vi = std::max(0,std::min(Vindex,nbins_v-1))

Definition at line 88 of file gpu_pitch_angle_diffusion.cpp.

◆ Vindex

const int Vindex = static_cast<int>(std::nearbyint(floor((normV) / dVbins)))

Definition at line 80 of file gpu_pitch_angle_diffusion.cpp.

◆ Vmu

const Real Vmu = dVbins * (Vindex+0.5)

Definition at line 81 of file gpu_pitch_angle_diffusion.cpp.

◆ Vpara

const Real Vpara = VplasmaX*dev_bValues[3*cellIdx] + VplasmaY*dev_bValues[3*cellIdx+1] + VplasmaZ*dev_bValues[3*cellIdx+2]

Definition at line 77 of file gpu_pitch_angle_diffusion.cpp.

◆ VplasmaX

const Real VplasmaX = VX - dev_bulkVX[cellIdx]

Definition at line 72 of file gpu_pitch_angle_diffusion.cpp.

◆ VplasmaY

const Real VplasmaY = VY - dev_bulkVY[cellIdx]

Definition at line 73 of file gpu_pitch_angle_diffusion.cpp.

◆ VplasmaZ

const Real VplasmaZ = VZ - dev_bulkVZ[cellIdx]

Definition at line 74 of file gpu_pitch_angle_diffusion.cpp.

◆ VX

const Real VX = blockParameters[BlockParams::VXCRD] + (i + 0.5)*blockParameters[BlockParams::DVX]

Definition at line 68 of file gpu_pitch_angle_diffusion.cpp.

◆ VY

const Real VY = blockParameters[BlockParams::VYCRD] + (j + 0.5)*blockParameters[BlockParams::DVY]

Definition at line 69 of file gpu_pitch_angle_diffusion.cpp.

◆ VZ

const Real VZ = blockParameters[BlockParams::VZCRD] + (k + 0.5)*blockParameters[BlockParams::DVZ]

Definition at line 70 of file gpu_pitch_angle_diffusion.cpp.