cwpearson/astaroth
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Removed all MPI-related code in preparation of a rewrite of the MPI stuff

Browse Source
This commit is contained in:
jpekkila
2020-01-17 14:22:11 +02:00
parent 9264b7515a
commit 975a15f7f4
3 changed files with 2 additions and 2410 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2274
src/core/device.cc
View File

File diff suppressed because it is too large Load Diff

120
src/core/kernels/packing.cu
View File

@@ -28,93 +28,6 @@
#include "src/core/errchk.h" #include "src/core/errchk.h"
/*
__global__ void
kernel_pack_data(const AcReal* unpacked, const int3 unpacked_start, const int3 packed_dimensions,
AcReal* packed)
{
const int i_packed = threadIdx.x + blockIdx.x * blockDim.x;
const int j_packed = threadIdx.y + blockIdx.y * blockDim.y;
const int k_packed = threadIdx.z + blockIdx.z * blockDim.z;
// If within the start-end range (this allows threadblock dims that are not
// divisible by end - start)
if (i_packed >= packed_dimensions.x || //
j_packed >= packed_dimensions.y || //
k_packed >= packed_dimensions.z) {
return;
}
const int i_unpacked = i_packed + unpacked_start.x;
const int j_unpacked = j_packed + unpacked_start.y;
const int k_unpacked = k_packed + unpacked_start.z;
const int unpacked_idx = DEVICE_VTXBUF_IDX(i_unpacked, j_unpacked, k_unpacked);
const int packed_idx = i_packed + //
j_packed * packed_dimensions.x + //
k_packed * packed_dimensions.x * packed_dimensions.y;
packed[packed_idx] = unpacked[unpacked_idx];
}
__global__ void
kernel_unpack_data(const AcReal* packed, const int3 packed_dimensions, const int3 unpacked_start,
AcReal* unpacked)
{
const int i_packed = threadIdx.x + blockIdx.x * blockDim.x;
const int j_packed = threadIdx.y + blockIdx.y * blockDim.y;
const int k_packed = threadIdx.z + blockIdx.z * blockDim.z;
// If within the start-end range (this allows threadblock dims that are not
// divisible by end - start)
if (i_packed >= packed_dimensions.x || //
j_packed >= packed_dimensions.y || //
k_packed >= packed_dimensions.z) {
return;
}
const int i_unpacked = i_packed + unpacked_start.x;
const int j_unpacked = j_packed + unpacked_start.y;
const int k_unpacked = k_packed + unpacked_start.z;
const int unpacked_idx = DEVICE_VTXBUF_IDX(i_unpacked, j_unpacked, k_unpacked);
const int packed_idx = i_packed + //
j_packed * packed_dimensions.x + //
k_packed * packed_dimensions.x * packed_dimensions.y;
unpacked[unpacked_idx] = packed[packed_idx];
}
AcResult
acKernelPackData(const cudaStream_t stream, const AcReal* unpacked, const int3 unpacked_start,
const int3 packed_dimensions, AcReal* packed)
{
const dim3 tpb(32, 8, 1);
const dim3 bpg((unsigned int)ceil(packed_dimensions.x / (float)tpb.x),
(unsigned int)ceil(packed_dimensions.y / (float)tpb.y),
(unsigned int)ceil(packed_dimensions.z / (float)tpb.z));
kernel_pack_data<<<bpg, tpb, 0, stream>>>(unpacked, unpacked_start, packed_dimensions, packed);
ERRCHK_CUDA_KERNEL_ALWAYS(); // TODO SET W/ DEBUG ONLY
return AC_SUCCESS;
}
AcResult
acKernelUnpackData(const cudaStream_t stream, const AcReal* packed, const int3 packed_dimensions,
const int3 unpacked_start, AcReal* unpacked)
{
const dim3 tpb(32, 8, 1);
const dim3 bpg((unsigned int)ceil(packed_dimensions.x / (float)tpb.x),
(unsigned int)ceil(packed_dimensions.y / (float)tpb.y),
(unsigned int)ceil(packed_dimensions.z / (float)tpb.z));
kernel_unpack_data<<<bpg, tpb, 0, stream>>>(packed, packed_dimensions, unpacked_start,
unpacked);
ERRCHK_CUDA_KERNEL_ALWAYS(); // TODO SET W/ DEBUG ONLY
return AC_SUCCESS;
}*/
__global__ void __global__ void
kernel_pack_data(const VertexBufferArray vba, const int3 vba_start, PackedData packed) kernel_pack_data(const VertexBufferArray vba, const int3 vba_start, PackedData packed)
{ {
@@ -201,36 +114,3 @@ acKernelUnpackData(const cudaStream_t stream, const PackedData packed, const int
ERRCHK_CUDA_KERNEL_ALWAYS(); // TODO SET W/ DEBUG ONLY ERRCHK_CUDA_KERNEL_ALWAYS(); // TODO SET W/ DEBUG ONLY
return AC_SUCCESS; return AC_SUCCESS;
} }
AcResult
acKernelPackCorner(void)
{
return AC_FAILURE;
}
AcResult
acKernelUnpackCorner(void)
{
return AC_FAILURE;
}
AcResult
acKernelPackEdge(void)
{
return AC_FAILURE;
}
AcResult
acKernelUnpackEdge(void)
{
return AC_FAILURE;
}
AcResult
acKernelPackSide(void)
{
return AC_FAILURE;
}
AcResult
acKernelUnpackSide(void)
{
return AC_FAILURE;
}

18
src/core/kernels/packing.cuh
View File

@@ -28,15 +28,6 @@
#include "astaroth.h" #include "astaroth.h"
#include "common.cuh" #include "common.cuh"
/*
AcResult acKernelPackData(const cudaStream_t stream, const AcReal* unpacked,
const int3 unpacked_start, const int3 packed_dimensions, AcReal* packed);
AcResult acKernelUnpackData(const cudaStream_t stream, const AcReal* packed,
const int3 packed_dimensions, const int3 unpacked_start,
AcReal* unpacked);
*/
typedef struct { typedef struct {
int3 dims; int3 dims;
AcReal* data; AcReal* data;
@@ -47,12 +38,3 @@ AcResult acKernelPackData(const cudaStream_t stream, const VertexBufferArray vba
AcResult acKernelUnpackData(const cudaStream_t stream, const PackedData packed, AcResult acKernelUnpackData(const cudaStream_t stream, const PackedData packed,
const int3 vba_start, VertexBufferArray vba); const int3 vba_start, VertexBufferArray vba);
AcResult acKernelPackCorner(void);
AcResult acKernelUnpackCorner(void);
AcResult acKernelPackEdge(void);
AcResult acKernelUnpackEdge(void);
AcResult acKernelPackSide(void);
AcResult acKernelUnpackSide(void);