Skip to content

Commit

Permalink
struct template for fhp_grid
Browse files Browse the repository at this point in the history
  • Loading branch information
loonatick-src committed May 11, 2021
1 parent d4777e3 commit 7ebe472
Show file tree
Hide file tree
Showing 5 changed files with 94 additions and 12 deletions.
10 changes: 10 additions & 0 deletions TODO
Original file line number Diff line number Diff line change
@@ -0,0 +1,10 @@
struct fhp_grid
-> velocities (device vector)
-> word * grid
-> __device__ functions
-> stream
-> collide
-> momentum
-> average momentum(radius)
-> particle_count
}
Binary file removed build/libfhp.a
Binary file not shown.
85 changes: 84 additions & 1 deletion src/fhp.hpp
Original file line number Diff line number Diff line change
@@ -1,7 +1,90 @@
# pragma once
#include "dbg.h"
#include <thrust/device_vector>
#include <thrust/host_vector>
#include "velocity.hpp"


typedef uint8_t u8;
typedef uint32_t u32;
typedef velocity<2, double> velocity2;

template <typename word, size_t BLOCK_WIDTH, size_t BLOCK_HEIGHT = BLOCK_WIDTH>
struct fhp_grid
{

// do we want to go all in and make device_grid thrust::device_vector?
// TOOD: initialize lookup tables (as members?)
word *device_grid;
const size_t width, height;
const thrust::device_vector<velocity2> channels; // will this be directly acessible inside the __global__ and __device__ functions?
#if 0
fhp_grid(size_t width, size_t height, thrust::host_vector<velocity2> velocities) :
width{width}, height{height}, channels {velocities};
#endif
fhp_grid(size_t w, size_t h, thrust::host_vector<velocity2> velocities, word *buffer) :
width {w}, height{h}, channels {velocities}
{
assert(buffer != NULL);
const auto grid_sz = width * height;
const auto mem_sz = grid_sz * sizeof(word);
cudaMalloc((void **) device_grid, mem_sz);
cudaMemcpy(device_grid, buffer, mem_sz);
}

~fhp_grid()
{
cudaFree(device_grid);
}

template <size_t timesteps>
void start_evolution()
{
// TODO: port from master to this
/*
assert(channels.size() > 0);
const dim3 block_config(BLOCK_HEIGHT, BLOCK_WIDTH);
const dim3 grid_config(...);
evolve<<<>>>();
*/
}

template <size_t timesteps>
__global__
void evolve()
{
/*
#pragma unroll
for (size_t t = 0; t < timesteps; t++)
{
}
*/

}

__device__
void stream()
{
// apply streaming operator to grid
}

__device__
void collide()
{
// apply collision operator to grid
}

__device__
void momentum()
{
// calculate momentum at threadId, blockId
}

__device__
void occupancy()
{
// calculate occupancy at threadId, blockId
// not to be confused with CUDA occupancy
}
};
11 changes: 0 additions & 11 deletions tests/tests.log

This file was deleted.

Binary file removed tests/velocity_tests
Binary file not shown.

0 comments on commit 7ebe472

Please sign in to comment.