Skip to content

Commit

Permalink
compilation typo fix, consolidate communication bitmasks into spatial…
Browse files Browse the repository at this point in the history 
…_cells/spatial_cell_wrapper.hpp
  • Loading branch information
@markusbattarbee
markusbattarbee committed Jan 30, 2025
1 parent c5110ab commit a8d86cf
Show file tree
Hide file tree
Showing 7 changed files with 90 additions and 148 deletions.
2 changes: 1 addition & 1 deletion arch/gpu_base.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -404,7 +404,7 @@ __host__ void gpu_vlasov_allocate(
uint maxBlockCount // Largest found vmesh size
) {
// Always prepare for at least VLASOV_BUFFER_MINBLOCKS blocks
const uint maxBlocksPerCell = max(VLASOV_BUFFER_MINBLOCKS, maxBlockCount)
const uint maxBlocksPerCell = max(VLASOV_BUFFER_MINBLOCKS, maxBlockCount);
const uint maxNThreads = gpu_getMaxThreads();
for (uint i=0; i<maxNThreads; ++i) {
gpu_vlasov_allocate_perthread(i, maxBlocksPerCell);
Expand Down
4 changes: 2 additions & 2 deletions arch/gpu_base.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -53,8 +53,8 @@
// #define INIT_VMESH_SIZE 4096
// #define INIT_MAP_SIZE 14 // 2^14 = 16384

static const double VLASOV_BUFFER_MINBLOCKS = 2500;
static const double VLASOV_BUFFER_MINCOLUMNS = 500;
static const uint VLASOV_BUFFER_MINBLOCKS = 2500;
static const uint VLASOV_BUFFER_MINCOLUMNS = 500;
static const double BLOCK_ALLOCATION_PADDING = 1.2;
static const double BLOCK_ALLOCATION_FACTOR = 1.1;
// static const double BLOCK_ALLOCATION_PADDING = 1.5;
Expand Down
2 changes: 1 addition & 1 deletion spatial_cells/spatial_cell_cpu.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,7 @@

#include <unordered_set>

#include "spatial_cell_cpu.hpp"
#include "spatial_cell_wrapper.hpp"
#include "../object_wrapper.h"

#ifdef DEBUG_VLASIATOR
Expand Down
81 changes: 10 additions & 71 deletions spatial_cells/spatial_cell_cpu.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -56,79 +56,8 @@ Spatial cell class for Vlasiator that supports a variable number of velocity blo
#endif
#endif

typedef Parameters P; // Needed in numerous files which include this one

/*!
Used as an error from functions returning velocity cells or
as a cell that would be outside of the velocity block
*/
#define error_velocity_cell 0xFFFFFFFFu

/*!
Used as an error from functions returning velocity cell indices or
as an index that would be outside of the velocity block
*/
#define error_velocity_cell_index 0xFFFFFFFFu

namespace spatial_cell {

namespace Transfer {
const uint64_t NONE = 0;
const uint64_t CELL_PARAMETERS = (1ull<<0);
const uint64_t CELL_DERIVATIVES = (1ull<<1);
const uint64_t VEL_BLOCK_LIST_STAGE1 = (1ull<<2);
const uint64_t VEL_BLOCK_LIST_STAGE2 = (1ull<<3);
const uint64_t VEL_BLOCK_DATA = (1ull<<4);
const uint64_t VEL_BLOCK_PARAMETERS = (1ull<<6);
const uint64_t VEL_BLOCK_WITH_CONTENT_STAGE1 = (1ull<<7);
const uint64_t VEL_BLOCK_WITH_CONTENT_STAGE2 = (1ull<<8);
const uint64_t CELL_SYSBOUNDARYFLAG = (1ull<<9);
const uint64_t CELL_E = (1ull<<10);
const uint64_t CELL_EDT2 = (1ull<<11);
const uint64_t CELL_PERB = (1ull<<12);
const uint64_t CELL_PERBDT2 = (1ull<<13);
const uint64_t CELL_RHOM_V = (1ull<<14);
const uint64_t CELL_RHOMDT2_VDT2 = (1ull<<15);
const uint64_t CELL_RHOQ = (1ull<<16);
const uint64_t CELL_RHOQDT2 = (1ull<<17);
const uint64_t CELL_BVOL = (1ull<<18);
const uint64_t CELL_BVOL_DERIVATIVES = (1ull<<19);
const uint64_t CELL_DIMENSIONS = (1ull<<20);
const uint64_t CELL_IOLOCALCELLID = (1ull<<21);
const uint64_t NEIGHBOR_VEL_BLOCK_DATA = (1ull<<22);
const uint64_t CELL_HALL_TERM = (1ull<<23);
const uint64_t CELL_P = (1ull<<24);
const uint64_t CELL_PDT2 = (1ull<<25);
const uint64_t POP_METADATA = (1ull<<26);
const uint64_t RANDOMGEN = (1ull<<27);
const uint64_t CELL_GRADPE_TERM = (1ull<<28);
const uint64_t REFINEMENT_PARAMETERS = (1ull<<29);
//all data
const uint64_t ALL_DATA =
CELL_PARAMETERS
| CELL_DERIVATIVES | CELL_BVOL_DERIVATIVES
| VEL_BLOCK_DATA
| CELL_SYSBOUNDARYFLAG
| POP_METADATA | RANDOMGEN;

//all data, except the distribution function
const uint64_t ALL_SPATIAL_DATA =
CELL_PARAMETERS
| CELL_DERIVATIVES | CELL_BVOL_DERIVATIVES
| CELL_SYSBOUNDARYFLAG
| POP_METADATA | RANDOMGEN;
}

typedef std::array<unsigned int, 3> velocity_cell_indices_t; /**< Defines the indices of a velocity cell in a velocity block.
* Indices start from 0 and the first value is the index in x direction.
* Note: these are the (i,j,k) indices of the cell within the block.
* Valid values are ([0,WID[,[0,WID[,[0,WID[).*/

typedef std::array<vmesh::LocalID,3> velocity_block_indices_t; /**< Defines the indices of a velocity block in the velocity grid.
* Indices start from 0 and the first value is the index in x direction.
* Note: these are the (i,j,k) indices of the block.
* Valid values are ([0,vx_length[,[0,vy_length[,[0,vz_length[).*/

/** Wrapper for variables needed for each particle species.
* Change order if you know what you are doing.
* All Real fields should be consecutive, as they are communicated as a block.
Expand Down Expand Up @@ -277,6 +206,16 @@ namespace spatial_cell {
}
};

typedef std::array<unsigned int, 3> velocity_cell_indices_t; /**< Defines the indices of a velocity cell in a velocity block.
* Indices start from 0 and the first value is the index in x direction.
* Note: these are the (i,j,k) indices of the cell within the block.
* Valid values are ([0,WID[,[0,WID[,[0,WID[).*/

typedef std::array<vmesh::LocalID,3> velocity_block_indices_t; /**< Defines the indices of a velocity block in the velocity grid.
* Indices start from 0 and the first value is the index in x direction.
* Note: these are the (i,j,k) indices of the block.
* Valid values are ([0,vx_length[,[0,vy_length[,[0,vz_length[).*/

class SpatialCell {
public:
SpatialCell();
Expand Down
2 changes: 1 addition & 1 deletion spatial_cells/spatial_cell_gpu.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,7 @@

#include <unordered_set>

#include "spatial_cell_gpu.hpp"
#include "spatial_cell_wrapper.hpp"
#include "../arch/gpu_base.hpp"
#include "../object_wrapper.h"
#include "../velocity_mesh_parameters.h"
Expand Down
82 changes: 10 additions & 72 deletions spatial_cells/spatial_cell_gpu.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -55,80 +55,8 @@ Spatial cell class for Vlasiator that supports a variable number of velocity blo
#endif
#endif

typedef Parameters P; // Heeded in numerous files which include this one

/*!
Used as an error from functions returning velocity cells or
as a cell that would be outside of the velocity block
*/
#define error_velocity_cell 0xFFFFFFFFu

/*!
Used as an error from functions returning velocity cell indices or
as an index that would be outside of the velocity block
*/
#define error_velocity_cell_index 0xFFFFFFFFu

namespace spatial_cell {

namespace Transfer {
const uint64_t NONE = 0;
const uint64_t CELL_PARAMETERS = (1ull<<0);
const uint64_t CELL_DERIVATIVES = (1ull<<1);
const uint64_t VEL_BLOCK_LIST_STAGE1 = (1ull<<2);
const uint64_t VEL_BLOCK_LIST_STAGE2 = (1ull<<3);
const uint64_t VEL_BLOCK_DATA = (1ull<<4);
const uint64_t VEL_BLOCK_PARAMETERS = (1ull<<6);
const uint64_t VEL_BLOCK_WITH_CONTENT_STAGE1 = (1ull<<7);
const uint64_t VEL_BLOCK_WITH_CONTENT_STAGE2 = (1ull<<8);
const uint64_t CELL_SYSBOUNDARYFLAG = (1ull<<9);
const uint64_t CELL_E = (1ull<<10);
const uint64_t CELL_EDT2 = (1ull<<11);
const uint64_t CELL_PERB = (1ull<<12);
const uint64_t CELL_PERBDT2 = (1ull<<13);
const uint64_t CELL_RHOM_V = (1ull<<14);
const uint64_t CELL_RHOMDT2_VDT2 = (1ull<<15);
const uint64_t CELL_RHOQ = (1ull<<16);
const uint64_t CELL_RHOQDT2 = (1ull<<17);
const uint64_t CELL_BVOL = (1ull<<18);
const uint64_t CELL_BVOL_DERIVATIVES = (1ull<<19);
const uint64_t CELL_DIMENSIONS = (1ull<<20);
const uint64_t CELL_IOLOCALCELLID = (1ull<<21);
const uint64_t NEIGHBOR_VEL_BLOCK_DATA = (1ull<<22);
const uint64_t CELL_HALL_TERM = (1ull<<23);
const uint64_t CELL_P = (1ull<<24);
const uint64_t CELL_PDT2 = (1ull<<25);
const uint64_t POP_METADATA = (1ull<<26);
const uint64_t RANDOMGEN = (1ull<<27);
const uint64_t CELL_GRADPE_TERM = (1ull<<28);
const uint64_t REFINEMENT_PARAMETERS = (1ull<<29);
//all data
const uint64_t ALL_DATA =
CELL_PARAMETERS
| CELL_DERIVATIVES | CELL_BVOL_DERIVATIVES
| VEL_BLOCK_DATA
| CELL_SYSBOUNDARYFLAG
| POP_METADATA | RANDOMGEN;

//all data, except the distribution function
const uint64_t ALL_SPATIAL_DATA =
CELL_PARAMETERS
| CELL_DERIVATIVES | CELL_BVOL_DERIVATIVES
| CELL_SYSBOUNDARYFLAG
| POP_METADATA | RANDOMGEN;
}

typedef std::array<unsigned int, 3> velocity_cell_indices_t; /**< Defines the indices of a velocity cell in a velocity block.
* Indices start from 0 and the first value is the index in x direction.
* Note: these are the (i,j,k) indices of the cell within the block.
* Valid values are ([0,WID[,[0,WID[,[0,WID[).*/

typedef std::array<vmesh::LocalID,3> velocity_block_indices_t; /**< Defines the indices of a velocity block in the velocity grid.
* Indices start from 0 and the first value is the index in x direction.
* Note: these are the (i,j,k) indices of the block.
* Valid values are ([0,vx_length[,[0,vy_length[,[0,vz_length[).*/


/** GPU mini-kernel for resizing a vmesh on-device */
__global__ static void resize_vmesh_ondevice_kernel (
vmesh::VelocityMesh *vmesh,
Expand Down Expand Up @@ -592,6 +520,16 @@ __global__ static void resize_and_empty_kernel (
}
}

typedef std::array<unsigned int, 3> velocity_cell_indices_t; /**< Defines the indices of a velocity cell in a velocity block.
* Indices start from 0 and the first value is the index in x direction.
* Note: these are the (i,j,k) indices of the cell within the block.
* Valid values are ([0,WID[,[0,WID[,[0,WID[).*/

typedef std::array<vmesh::LocalID,3> velocity_block_indices_t; /**< Defines the indices of a velocity block in the velocity grid.
* Indices start from 0 and the first value is the index in x direction.
* Note: these are the (i,j,k) indices of the block.
* Valid values are ([0,vx_length[,[0,vy_length[,[0,vz_length[).*/

class SpatialCell {
public:
SpatialCell();
Expand Down
65 changes: 65 additions & 0 deletions spatial_cells/spatial_cell_wrapper.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -31,4 +31,69 @@
#include "spatial_cell_cpu.hpp"
#endif

typedef Parameters P; // Heeded in numerous files which include this one

/*!
Used as an error from functions returning velocity cells or
as a cell that would be outside of the velocity block
*/
#define error_velocity_cell 0xFFFFFFFFu

/*!
Used as an error from functions returning velocity cell indices or
as an index that would be outside of the velocity block
*/
#define error_velocity_cell_index 0xFFFFFFFFu

namespace spatial_cell {

namespace Transfer {
static const uint64_t NONE = 0;
static const uint64_t CELL_PARAMETERS = (1ull<<0);
static const uint64_t CELL_DERIVATIVES = (1ull<<1);
static const uint64_t VEL_BLOCK_LIST_STAGE1 = (1ull<<2);
static const uint64_t VEL_BLOCK_LIST_STAGE2 = (1ull<<3);
static const uint64_t VEL_BLOCK_DATA = (1ull<<4);
static const uint64_t VEL_BLOCK_PARAMETERS = (1ull<<6);
static const uint64_t VEL_BLOCK_WITH_CONTENT_STAGE1 = (1ull<<7);
static const uint64_t VEL_BLOCK_WITH_CONTENT_STAGE2 = (1ull<<8);
static const uint64_t CELL_SYSBOUNDARYFLAG = (1ull<<9);
static const uint64_t CELL_E = (1ull<<10);
static const uint64_t CELL_EDT2 = (1ull<<11);
static const uint64_t CELL_PERB = (1ull<<12);
static const uint64_t CELL_PERBDT2 = (1ull<<13);
static const uint64_t CELL_RHOM_V = (1ull<<14);
static const uint64_t CELL_RHOMDT2_VDT2 = (1ull<<15);
static const uint64_t CELL_RHOQ = (1ull<<16);
static const uint64_t CELL_RHOQDT2 = (1ull<<17);
static const uint64_t CELL_BVOL = (1ull<<18);
static const uint64_t CELL_BVOL_DERIVATIVES = (1ull<<19);
static const uint64_t CELL_DIMENSIONS = (1ull<<20);
static const uint64_t CELL_IOLOCALCELLID = (1ull<<21);
static const uint64_t NEIGHBOR_VEL_BLOCK_DATA = (1ull<<22);
static const uint64_t CELL_HALL_TERM = (1ull<<23);
static const uint64_t CELL_P = (1ull<<24);
static const uint64_t CELL_PDT2 = (1ull<<25);
static const uint64_t POP_METADATA = (1ull<<26);
static const uint64_t RANDOMGEN = (1ull<<27);
static const uint64_t CELL_GRADPE_TERM = (1ull<<28);
static const uint64_t REFINEMENT_PARAMETERS = (1ull<<29);
//all data
static const uint64_t ALL_DATA =
CELL_PARAMETERS
| CELL_DERIVATIVES | CELL_BVOL_DERIVATIVES
| VEL_BLOCK_DATA
| CELL_SYSBOUNDARYFLAG
| POP_METADATA | RANDOMGEN;

//all data, except the distribution function
static const uint64_t ALL_SPATIAL_DATA =
CELL_PARAMETERS
| CELL_DERIVATIVES | CELL_BVOL_DERIVATIVES
| CELL_SYSBOUNDARYFLAG
| POP_METADATA | RANDOMGEN;
}

}

#endif

0 comments on commit a8d86cf

Please sign in to comment.