Skip to content

Commit

Permalink
Sorted out the macro for conditionally accessing OneAPI extensions
Browse files Browse the repository at this point in the history
Added idx() to the non Sycl CellBase class
Added OpenMP pragma around forALl loop
Implemented pattern of passing pointers as kernel parameters for stack defined variables
Implemented an id attribute on the EcolabCell class
  • Loading branch information
highperformancecoder committed Dec 11, 2024
1 parent d8de4f4 commit 7167d78
Show file tree
Hide file tree
Showing 5 changed files with 45 additions and 27 deletions.
17 changes: 13 additions & 4 deletions include/ecolab.h
Original file line number Diff line number Diff line change
Expand Up @@ -20,11 +20,10 @@
#include "device/Ouroboros_impl.dp.hpp"
#include "InstanceDefinitions.dp.hpp"
#include "device/MemoryInitialization.dp.hpp"
#if 1 //def __INTEL_LLVM_COMPLER
#warning "Engaging experimental printf"
#ifdef __INTEL_LLVM_COMPILER
template <typename... Args>
void syclPrintf(Args... args) {sycl::ext::oneapi::experimental::printf(args...);}
sycl::nd_item<1> syclItem() {return sycl::ext::oneapi::this_work_item::get_nd_item<1>();}
inline sycl::nd_item<1> syclItem() {return sycl::ext::oneapi::this_work_item::get_nd_item<1>();}
#else
#define syclPrintf(...)
#endif
Expand Down Expand Up @@ -176,6 +175,8 @@ namespace ecolab
#else
template <class T> using CellAllocator=std::allocator<T>;
template <class T> CellAllocator<T> allocator() const {return CellAllocator<T>();}
size_t m_idx=0; // stash the position within the local node vector here
size_t idx() const {return m_idx;}
#endif
};

Expand Down Expand Up @@ -247,7 +248,15 @@ namespace ecolab
});
syclQ().wait();
#else
for (auto& i: *this) f(*i->template as<Cell>());
auto sz=this->size();
#ifdef _OPENMP
#pragma omp parallel for
#endif
for (size_t idx=0; idx<sz; ++idx) {
auto& cell=*(*this)[idx]->template as<Cell>();
cell.m_idx=idx;
f(cell);
}
#endif
}
};
Expand Down
2 changes: 0 additions & 2 deletions include/sparse_mat.h
Original file line number Diff line number Diff line change
Expand Up @@ -14,8 +14,6 @@

#include <arrays.h>

using sycl::ext::oneapi::experimental::printf;

namespace ecolab
{
/// sparse matrix class
Expand Down
31 changes: 17 additions & 14 deletions models/ecolab_model.cc
Original file line number Diff line number Diff line change
Expand Up @@ -185,19 +185,20 @@ void SpatialModel::mutate()
{
last_mut_tstep=tstep;

DeviceType<array<Float,graphcode::Allocator<Float>>> mut_scale;
#ifdef SYCL_LANGUAGE_VERSION
using ArrayAlloc=CellBase::CellAllocator<unsigned>;
using NewSpAlloc=graphcode::Allocator<array<unsigned,ArrayAlloc>>;
vector<array<unsigned,ArrayAlloc>,NewSpAlloc> newSpV
vector<array<unsigned,ArrayAlloc>,NewSpAlloc> newSp
(size(),NewSpAlloc(syclQ(),sycl::usm::alloc::shared));
auto newSp=newSpV.data();
mut_scale->allocator(graphcode::Allocator<Float>(syclQ(),sycl::usm::alloc::shared));
#else
vector<array<unsigned>> newSp(size());
#endif
*mut_scale=sp_sep * repro_rate * mutation * int(tstep - last_mut_tstep);

forAll([=,this](EcolabCell& c) {
auto mut_scale=sp_sep * repro_rate * mutation * int(tstep - last_mut_tstep);
newSp[c.idx()] = c.mutate(mut_scale);
forAll([=,newSp=newSp.data(),mut_scale=&*mut_scale,this](EcolabCell& c) {
newSp[c.idx()] = c.mutate(*mut_scale);
});

array<unsigned> new_sp;
Expand Down Expand Up @@ -233,9 +234,8 @@ void SpatialModel::mutate()
ModelData::mutate(new_sp);
#endif
// set the new species density to 1 for those created on this cell
auto cell_ids_p=&*cell_ids;
forAll([=,this](EcolabCell& c) {
c.density <<= (*cell_ids_p)==(*this)[c.idx()].id();
forAll([=,cell_ids=&*cell_ids,this](EcolabCell& c) {
c.density <<= (*cell_ids)==(*this)[c.idx()].id();
});
}

Expand Down Expand Up @@ -494,7 +494,9 @@ void SpatialModel::setGrid(size_t nx, size_t ny)
for (size_t j=0; j<numY; ++j)
{
auto o=insertObject(makeId(i,j));
o->as<EcolabCell>()->rand.seed(o.id());
auto& c=*o->as<EcolabCell>();
c.id=o.id();
c.rand.seed(o.id());
o.proc(o.id() % nprocs()); // TODO can we get this to work without this.
// wire up von Neumann neighborhood
o->neighbours.push_back(makeId(i-1,j));
Expand Down Expand Up @@ -526,30 +528,31 @@ void SpatialModel::migrate()
using Array=vector<int,ArrayAlloc>;
using DeltaAlloc=graphcode::Allocator<Array>;
Array init(species.size(),0,ArrayAlloc(syclQ(),sycl::usm::alloc::device));
vector<Array,DeltaAlloc> deltaV(size(), init, DeltaAlloc(syclQ(),sycl::usm::alloc::device));
auto delta=deltaV.data();
vector<Array,DeltaAlloc> delta(size(), init, DeltaAlloc(syclQ(),sycl::usm::alloc::device));
#else
vector<array<int>> delta(size(), array<int>(species.size(),0));
#endif

forAll([=,this](EcolabCell& c) {
forAll([=,delta=delta.data(),this](EcolabCell& c) {
using FArray=array<Float,EcolabCell::CellAllocator<Float>>;
/* loop over neighbours */
for (auto& n: c)
{
auto& nbr=*n->as<EcolabCell>();
FArray m( Float(tstep-last_mig_tstep) * migration *
(nbr.density - c.density), c.allocator<Float>());
Float salt=(*this)[c.idx()].id()<n.id()? c.salt: nbr.salt;
Float salt=c.id<nbr.id? c.salt: nbr.salt;
// array::operator+= cannot be used in kernels
// delta[c.idx()]+=m + (m!=0.0)*(2*(m>0.0)-1)) * salt;
FArray tmp=(m + (m!=0.0)*(2*(m>0.0)-1)) * salt;
array_ns::asg_plus_v(delta[c.idx()].data(),m.size(), tmp);
}
});
last_mig_tstep=tstep;
#ifdef SYCL_LANGUAGE_VERSION
syclQ().wait();
forAll([=,this](EcolabCell& c) {
#endif
forAll([=,delta=delta.data(),this](EcolabCell& c) {
// array::operator+= cannot be used in kernels
//c.density+=delta[c.idx()];
array_ns::asg_plus_v(c.density.data(), c.density.size(), delta[c.idx()]);
Expand Down
7 changes: 5 additions & 2 deletions models/ecolab_model.h
Original file line number Diff line number Diff line change
Expand Up @@ -60,7 +60,7 @@ template <class CellBase>
struct EcolabPoint: public Exclude<CellBase>
{
Float salt; /* random no. used for migration */
template <class T> using Allocator=CellBase::template CellAllocator<T>;
template <class T> using Allocator=typename CellBase::template CellAllocator<T>;
array<int,Allocator<int>> density{this->template allocator<int>()};
void generate(unsigned niter, const ModelData&);
void condense(const array<bool>& mask, size_t mask_true);
Expand Down Expand Up @@ -93,7 +93,10 @@ struct PanmicticModel: public ModelData, public EcolabPoint<AllocatorBase>, publ
array<double> lifetimes();
};

struct EcolabCell: public EcolabPoint<ecolab::CellBase>, public graphcode::Object<EcolabCell> {};
struct EcolabCell: public EcolabPoint<ecolab::CellBase>, public graphcode::Object<EcolabCell>
{
unsigned id=0; // stash the graphcode node id here
};

class SpatialModel: public ModelData, public EcolabGraph<EcolabCell>,
public ecolab::Model<SpatialModel>
Expand Down
15 changes: 10 additions & 5 deletions models/spatial_ecolab.py
Original file line number Diff line number Diff line change
Expand Up @@ -26,8 +26,8 @@ def randomList(num, min, max):

ecolab.species(range(nsp))

numX=8
numY=8
numX=32
numY=32
ecolab.setGrid(numX,numY)
ecolab.partitionObjects()

Expand All @@ -52,12 +52,17 @@ def randomList(num, min, max):

print(device())

def step():
def stepImpl():
ecolab.generate(100)
ecolab.mutate()
# ecolab.migrate()
# ecolab.condense()
ecolab.gather()

from timeit import timeit
print(timeit('stepImpl()', globals=globals(), number=10))

def step():
if myid()==0:
nsp=len(ecolab.species)
statusBar.configure(text=f't={ecolab.tstep()} nsp:{nsp}')
Expand All @@ -66,8 +71,8 @@ def step():
for i in range(numX):
for j in range(numY):
plot(f'Density({i},{j})',ecolab.tstep(),ecolab.cell(i,j).density(), pens=ecolab.species())
gui(step)

#gui(step)



0 comments on commit 7167d78

Please sign in to comment.