From 4ec94c97879aafef15f7663135745e4ba61e62cf Mon Sep 17 00:00:00 2001 From: Adrian Kummerlaender Date: Mon, 17 May 2021 00:15:33 +0200 Subject: Extract first public LiterateLB version --- tangle/LLBM/descriptor.h | 298 +++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 298 insertions(+) create mode 100644 tangle/LLBM/descriptor.h (limited to 'tangle/LLBM/descriptor.h') diff --git a/tangle/LLBM/descriptor.h b/tangle/LLBM/descriptor.h new file mode 100644 index 0000000..d4ca44d --- /dev/null +++ b/tangle/LLBM/descriptor.h @@ -0,0 +1,298 @@ +#pragma once + +#include +#include +#include +#include + +#ifdef __CUDA_ARCH__ + #define DATA device_data +#else + #define DATA host_data +#endif + +using pop_index_t = std::uint8_t; + +namespace descriptor { + +struct D2Q9 { + static constexpr unsigned d = 2; + static constexpr unsigned q = 9; +}; + +struct D3Q19 { + static constexpr unsigned d = 3; + static constexpr unsigned q = 19; +}; + +namespace device_data { + template + __constant__ pop_index_t opposite[DESCRIPTOR::q] { }; + + template + __constant__ int c[DESCRIPTOR::q][DESCRIPTOR::d] { }; + + template + __constant__ float c_length[DESCRIPTOR::q] { }; + + template + __constant__ float weight[DESCRIPTOR::q] { }; + + template <> + __constant__ pop_index_t opposite[9] = { + 8, 7, 6, 5, 4, 3, 2, 1, 0 + }; + + template <> + __constant__ int c[9][2] = { + {-1,-1}, {-1,0}, {-1,1}, {0,-1}, {0,0}, {0,1}, {1,-1}, {1,0}, {1,1} + }; + + template <> + __constant__ float c_length[9] = { + 1.41421356237310, 1.00000000000000, 1.41421356237310, 1.00000000000000, 0, 1.00000000000000, 1.41421356237310, 1.00000000000000, 1.41421356237310 + }; + + template <> + __constant__ float weight[9] = { + 0.0277777777777778, 0.111111111111111, 0.0277777777777778, 0.111111111111111, 0.444444444444444, 0.111111111111111, 0.0277777777777778, 0.111111111111111, 0.0277777777777778 + }; + + template <> + __constant__ pop_index_t opposite[19] = { + 18, 17, 16, 15, 14, 13, 12, 11, 10, 9, 8, 7, 6, 5, 4, 3, 2, 1, 0 + }; + + template <> + __constant__ int c[19][3] = { + {0,1,1}, {-1,0,1}, {0,0,1}, {1,0,1}, {0,-1,1}, {-1,1,0}, {0,1,0}, {1,1,0}, {-1,0,0}, {0,0,0}, {1,0,0}, {-1,-1,0}, {0,-1,0}, {1,-1,0}, {0,1,-1}, {-1,0,-1}, {0,0,-1}, {1,0,-1}, {0,-1,-1} + }; + + template <> + __constant__ float c_length[19] = { + 1.41421356237310, 1.41421356237310, 1.00000000000000, 1.41421356237310, 1.41421356237310, 1.41421356237310, 1.00000000000000, 1.41421356237310, 1.00000000000000, 0, 1.00000000000000, 1.41421356237310, 1.00000000000000, 1.41421356237310, 1.41421356237310, 1.41421356237310, 1.00000000000000, 1.41421356237310, 1.41421356237310 + }; + + template <> + __constant__ float weight[19] = { + 0.0277777777777778, 0.0277777777777778, 0.0555555555555556, 0.0277777777777778, 0.0277777777777778, 0.0277777777777778, 0.0555555555555556, 0.0277777777777778, 0.0555555555555556, 0.333333333333333, 0.0555555555555556, 0.0277777777777778, 0.0555555555555556, 0.0277777777777778, 0.0277777777777778, 0.0277777777777778, 0.0555555555555556, 0.0277777777777778, 0.0277777777777778 + }; +} + +namespace host_data { + template + constexpr pop_index_t opposite[DESCRIPTOR::q] { }; + + template + constexpr int c[DESCRIPTOR::q][DESCRIPTOR::d] { }; + + template + constexpr float c_length[DESCRIPTOR::q] { }; + + template + constexpr float weight[DESCRIPTOR::q] { }; + + template <> + constexpr pop_index_t opposite[9] = { + 8, 7, 6, 5, 4, 3, 2, 1, 0 + }; + + template <> + constexpr int c[9][2] = { + {-1,-1}, {-1,0}, {-1,1}, {0,-1}, {0,0}, {0,1}, {1,-1}, {1,0}, {1,1} + }; + + template <> + constexpr float c_length[9] = { + 1.41421356237310, 1.00000000000000, 1.41421356237310, 1.00000000000000, 0, 1.00000000000000, 1.41421356237310, 1.00000000000000, 1.41421356237310 + }; + + template <> + constexpr float weight[9] = { + 0.0277777777777778, 0.111111111111111, 0.0277777777777778, 0.111111111111111, 0.444444444444444, 0.111111111111111, 0.0277777777777778, 0.111111111111111, 0.0277777777777778 + }; + + template <> + constexpr pop_index_t opposite[19] = { + 18, 17, 16, 15, 14, 13, 12, 11, 10, 9, 8, 7, 6, 5, 4, 3, 2, 1, 0 + }; + + template <> + constexpr int c[19][3] = { + {0,1,1}, {-1,0,1}, {0,0,1}, {1,0,1}, {0,-1,1}, {-1,1,0}, {0,1,0}, {1,1,0}, {-1,0,0}, {0,0,0}, {1,0,0}, {-1,-1,0}, {0,-1,0}, {1,-1,0}, {0,1,-1}, {-1,0,-1}, {0,0,-1}, {1,0,-1}, {0,-1,-1} + }; + + template <> + constexpr float c_length[19] = { + 1.41421356237310, 1.41421356237310, 1.00000000000000, 1.41421356237310, 1.41421356237310, 1.41421356237310, 1.00000000000000, 1.41421356237310, 1.00000000000000, 0, 1.00000000000000, 1.41421356237310, 1.00000000000000, 1.41421356237310, 1.41421356237310, 1.41421356237310, 1.00000000000000, 1.41421356237310, 1.41421356237310 + }; + + template <> + constexpr float weight[19] = { + 0.0277777777777778, 0.0277777777777778, 0.0555555555555556, 0.0277777777777778, 0.0277777777777778, 0.0277777777777778, 0.0555555555555556, 0.0277777777777778, 0.0555555555555556, 0.333333333333333, 0.0555555555555556, 0.0277777777777778, 0.0555555555555556, 0.0277777777777778, 0.0277777777777778, 0.0277777777777778, 0.0555555555555556, 0.0277777777777778, 0.0277777777777778 + }; +} + +template +__host__ __device__ +pop_index_t opposite(pop_index_t iPop) { + return DESCRIPTOR::q - 1 - iPop; +} + +template +__host__ __device__ +int velocity(pop_index_t iPop, unsigned iDim) { + return DATA::template c[iPop][iDim]; +} + +template +__host__ __device__ +std::enable_if_t velocity(pop_index_t iPop) { + return make_float2(DATA::template c[iPop][0], + DATA::template c[iPop][1]); +} + +template +__host__ __device__ +std::enable_if_t velocity(pop_index_t iPop) { + return make_float3(DATA::template c[iPop][0], + DATA::template c[iPop][1], + DATA::template c[iPop][2]); +} + +template +__host__ __device__ +float velocity_length(pop_index_t iPop) { + return DATA::template c_length[iPop]; +} + +template +__host__ __device__ +float weight(pop_index_t iPop) { + return DATA::template weight[iPop]; +} + +template +struct CuboidD; + +template <> +struct CuboidD<2> { + const std::size_t nX; + const std::size_t nY; + const std::size_t nZ; + const std::size_t volume; + + CuboidD(std::size_t x, std::size_t y): + nX(x), nY(y), nZ(1), + volume(x*y) { }; +}; + +template <> +struct CuboidD<3> { + const std::size_t nX; + const std::size_t nY; + const std::size_t nZ; + const std::size_t volume; + const std::size_t plane; + + CuboidD(std::size_t x, std::size_t y, std::size_t z): + nX(x), nY(y), nZ(z), + volume(x*y*z), + plane(x*y) { }; +}; + +template +using Cuboid = CuboidD; + +__host__ __device__ +std::size_t gid(const CuboidD<2>& c, int iX, int iY, int iZ=0) { + return iY*c.nX + iX; +} + +__host__ __device__ +std::size_t gid(const CuboidD<3>& c, int iX, int iY, int iZ) { + return iZ*c.plane + iY*c.nX + iX; +} + +__host__ __device__ +int offset(const CuboidD<2>& c, int iX, int iY) { + return iY*c.nX + iX; +} + +template +__host__ __device__ +int offset(const CuboidD<2>& c, pop_index_t iPop) { + static_assert(DESCRIPTOR::d == 2, "Dimensions must match"); + return offset(c, + descriptor::velocity(iPop, 0), + descriptor::velocity(iPop, 1) + ); +} + +__host__ __device__ +int offset(const CuboidD<3>& c, int iX, int iY, int iZ) { + return iZ*c.plane + iY*c.nX + iX; +} + +template +__host__ __device__ +int offset(const CuboidD<3>& c, pop_index_t iPop) { + static_assert(DESCRIPTOR::d == 3, "Dimensions must match"); + return offset(c, + descriptor::velocity(iPop, 0), + descriptor::velocity(iPop, 1), + descriptor::velocity(iPop, 2) + ); +} + +template +__host__ __device__ +std::size_t neighbor(const CuboidD<2>& c, std::size_t iCell, pop_index_t iPop) { + return iCell + offset(c, iPop); +} + +template +__host__ __device__ +std::size_t neighbor(const CuboidD<3>& c, std::size_t iCell, pop_index_t iPop) { + return iCell + offset(c, iPop); +} + +__host__ __device__ +uint2 gidInverse(const CuboidD<2>& c, std::size_t gid) { + int iY = gid / c.nX; + int iX = gid % c.nX; + return make_uint2(iX, iY); +} + +__host__ __device__ +uint3 gidInverse(const CuboidD<3>& c, std::size_t gid) { + int iZ = gid / c.plane; + int iY = (gid % c.plane) / c.nX; + int iX = (gid % c.plane) % c.nX; + return make_uint3(iX,iY,iZ); +} + +__host__ __device__ +float2 gidInverseSmooth(const CuboidD<2>& c, std::size_t gid) { + int iY = gid / c.nX; + int iX = gid % c.nX; + return make_float2(iX, iY); +} + +__host__ __device__ +float3 gidInverseSmooth(const CuboidD<3>& c, std::size_t gid) { + int iZ = gid / c.plane; + int iY = (gid % c.plane) / c.nX; + int iX = (gid % c.plane) % c.nX; + return make_float3(iX,iY,iZ); +} + +bool isInside(const CuboidD<2>& c, std::size_t gid) { + return gid < c.volume; +} + +bool isInside(const CuboidD<3>& c, std::size_t gid) { + return gid < c.volume; +} + +} -- cgit v1.2.3