From 8bbb2cd5dfa9753c28d87d584d37b2340839a110 Mon Sep 17 00:00:00 2001 From: jpekkila Date: Mon, 12 Aug 2019 09:46:37 +0300 Subject: [PATCH 1/7] Now prints device info before trying to run the dummy kernel --- src/core/device.cu | 1 + 1 file changed, 1 insertion(+) diff --git a/src/core/device.cu b/src/core/device.cu index c81d468..89bb6b1 100644 --- a/src/core/device.cu +++ b/src/core/device.cu @@ -92,6 +92,7 @@ acDeviceCreate(const int id, const AcMeshInfo device_config, Device* device_hand device->id = id; device->local_config = device_config; + acDevicePrintInfo(device); // Check that the code was compiled for the proper GPU architecture printf("Trying to run a dummy kernel. If this fails, make sure that your\n" From b5daf22c260b34a17a4063a56f77c704c7f86d02 Mon Sep 17 00:00:00 2001 From: jpekkila Date: Mon, 12 Aug 2019 10:25:05 +0300 Subject: [PATCH 2/7] Added interface function acSynchronizeMesh --- include/astaroth.h | 3 +++ src/core/astaroth.cu | 6 ++++++ 2 files changed, 9 insertions(+) diff --git a/include/astaroth.h b/include/astaroth.h index cae7139..748c057 100644 --- a/include/astaroth.h +++ b/include/astaroth.h @@ -57,6 +57,9 @@ AcResult acCheckDeviceAvailability(void); * parameter*/ AcResult acSynchronizeStream(const Stream stream); +/** */ +AcResult acSynchronizeMesh(void); + /** Loads a constant to the memories of the devices visible to the caller */ AcResult acLoadDeviceConstant(const AcRealParam param, const AcReal value); diff --git a/src/core/astaroth.cu b/src/core/astaroth.cu index 6e06663..287bcd2 100644 --- a/src/core/astaroth.cu +++ b/src/core/astaroth.cu @@ -151,3 +151,9 @@ acLoadWithOffset(const AcMesh host_mesh, const int3 src, const int num_vertices) { return acNodeLoadMeshWithOffset(nodes[0], STREAM_DEFAULT, host_mesh, src, src, num_vertices); } + +AcResult +acSynchronizeMesh(void) +{ + return acNodeSynchronizeMesh(nodes[0], STREAM_DEFAULT); +} From bba9ec7c3b8302ffa37850925886040a569f7872 Mon Sep 17 00:00:00 2001 From: jpekkila Date: Mon, 12 Aug 2019 11:40:38 +0300 Subject: [PATCH 3/7] Implemented acNodeQueryDeviceConfiguration --- include/astaroth_node.h | 9 +++++++++ src/core/node.cu | 15 ++++++--------- 2 files changed, 15 insertions(+), 9 deletions(-) diff --git a/include/astaroth_node.h b/include/astaroth_node.h index 6b2d626..5458ac9 100644 --- a/include/astaroth_node.h +++ b/include/astaroth_node.h @@ -27,7 +27,16 @@ extern "C" { typedef struct node_s* Node; // Opaque pointer to node_s. typedef struct { + int3 m; + int3 n; +} Grid; +typedef struct { + int num_devices; + Device* devices; + + Grid grid; + Grid subgrid; } DeviceConfiguration; /** */ diff --git a/src/core/node.cu b/src/core/node.cu index f9a4aa9..6fa80f2 100644 --- a/src/core/node.cu +++ b/src/core/node.cu @@ -131,11 +131,6 @@ static const int MAX_NUM_DEVICES = 32; -typedef struct { - int3 m; - int3 n; -} Grid; - struct node_s { int id; @@ -334,10 +329,12 @@ acNodePrintInfo(const Node node) AcResult acNodeQueryDeviceConfiguration(const Node node, DeviceConfiguration* config) { - (void)node; - (void)config; - WARNING("Not implemented"); - return AC_FAILURE; + config->num_devices = node->num_devices; + config->devices = node->devices; + config->grid = node->grid; + config->subgrid = node->subgrid; + + return AC_SUCCESS; } AcResult From 3369d8efecd3d4203a70775a44266667bdc7b89e Mon Sep 17 00:00:00 2001 From: jpekkila Date: Mon, 12 Aug 2019 11:44:27 +0300 Subject: [PATCH 4/7] Added a missing include --- include/astaroth_node.h | 2 ++ 1 file changed, 2 insertions(+) diff --git a/include/astaroth_node.h b/include/astaroth_node.h index 5458ac9..61438f3 100644 --- a/include/astaroth_node.h +++ b/include/astaroth_node.h @@ -24,6 +24,8 @@ extern "C" { #include "astaroth_defines.h" +#include "astaroth_device.h" // TODO: Should this really be here? + typedef struct node_s* Node; // Opaque pointer to node_s. typedef struct { From e027f7e548bacd9d0f23160f7c70938733d28e03 Mon Sep 17 00:00:00 2001 From: jpekkila Date: Mon, 12 Aug 2019 13:25:47 +0300 Subject: [PATCH 5/7] Removed grid_n in astaroth.cu and replaced it with the new acNodeQueryDeviceConfiguration call --- src/core/astaroth.cu | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/src/core/astaroth.cu b/src/core/astaroth.cu index 287bcd2..fb7ab54 100644 --- a/src/core/astaroth.cu +++ b/src/core/astaroth.cu @@ -36,13 +36,10 @@ const char* vtxbuf_names[] = {AC_FOR_VTXBUF_HANDLES(AC_GEN_STR)}; static const int num_nodes = 1; static Node nodes[num_nodes]; -static int3 grid_n; AcResult acInit(const AcMeshInfo mesh_info) { - grid_n = (int3){mesh_info.int_params[AC_nx], mesh_info.int_params[AC_ny], - mesh_info.int_params[AC_nz]}; return acNodeCreate(0, mesh_info, &nodes[0]); } @@ -106,8 +103,11 @@ acIntegrate(const AcReal dt) AcResult acIntegrateStep(const int isubstep, const AcReal dt) { + DeviceConfiguration config; + acNodeQueryDeviceConfiguration(nodes[0], &config); + const int3 start = (int3){NGHOST, NGHOST, NGHOST}; - const int3 end = start + grid_n; + const int3 end = start + config.grid.n; return acNodeIntegrateSubstep(nodes[0], STREAM_DEFAULT, isubstep, start, end, dt); } From b8c4d07de259cf8644a27cda9a1e17000aa8c160 Mon Sep 17 00:00:00 2001 From: jpekkila Date: Mon, 12 Aug 2019 13:31:24 +0300 Subject: [PATCH 6/7] Removed unnecessary comments --- src/core/astaroth.cu | 5 ----- 1 file changed, 5 deletions(-) diff --git a/src/core/astaroth.cu b/src/core/astaroth.cu index fb7ab54..ea4ecaa 100644 --- a/src/core/astaroth.cu +++ b/src/core/astaroth.cu @@ -16,7 +16,6 @@ You should have received a copy of the GNU General Public License along with Astaroth. If not, see . */ -// #include "astaroth_defines.h" #include "astaroth.h" #include "errchk.h" @@ -93,10 +92,6 @@ acStore(AcMesh* host_mesh) AcResult acIntegrate(const AcReal dt) { - /* - acNodeIntegrate(nodes[0], dt); - return acBoundcondStep(); - */ return acNodeIntegrate(nodes[0], dt); } From d5b2e5bb42f42d405e1ab5ab8cd3454ee3316a32 Mon Sep 17 00:00:00 2001 From: jpekkila Date: Mon, 12 Aug 2019 14:05:35 +0300 Subject: [PATCH 7/7] Added placeholders for new built-in variables in the DSL. Also overloads to DCONST_INT etc. Naming still pending and old DCONST_REAL etc calls still work. --- src/core/device.cu | 38 ++++++++++++++++++++++++++++++++++---- 1 file changed, 34 insertions(+), 4 deletions(-) diff --git a/src/core/device.cu b/src/core/device.cu index 89bb6b1..9690278 100644 --- a/src/core/device.cu +++ b/src/core/device.cu @@ -40,14 +40,44 @@ typedef struct { } VertexBufferArray; __constant__ AcMeshInfo d_mesh_info; -#define DCONST_INT(X) (d_mesh_info.int_params[X]) -#define DCONST_INT3(X) (d_mesh_info.int3_params[X]) -#define DCONST_REAL(X) (d_mesh_info.real_params[X]) -#define DCONST_REAL3(X) (d_mesh_info.real3_params[X]) +static inline int __device__ +DCONST(const AcIntParam param) +{ + return d_mesh_info.int_params[param]; +} +static inline int3 __device__ +DCONST(const AcInt3Param param) +{ + return d_mesh_info.int3_params[param]; +} +static inline AcReal __device__ +DCONST(const AcRealParam param) +{ + return d_mesh_info.real_params[param]; +} +static inline AcReal3 __device__ +DCONST(const AcReal3Param param) +{ + return d_mesh_info.real3_params[param]; +} +#define DCONST_INT(x) DCONST(x) +#define DCONST_INT3(x) DCONST(x) +#define DCONST_REAL(x) DCONST(x) +#define DCONST_REAL3(x) DCONST(x) +//#define DCONST_INT(X) (d_mesh_info.int_params[X]) +//#define DCONST_INT3(X) (d_mesh_info.int3_params[X]) +//#define DCONST_REAL(X) (d_mesh_info.real_params[X]) +//#define DCONST_REAL3(X) (d_mesh_info.real3_params[X]) #define DEVICE_VTXBUF_IDX(i, j, k) ((i) + (j)*DCONST_INT(AC_mx) + (k)*DCONST_INT(AC_mxy)) #define DEVICE_1D_COMPDOMAIN_IDX(i, j, k) ((i) + (j)*DCONST_INT(AC_nx) + (k)*DCONST_INT(AC_nxy)) #define globalGridN (d_mesh_info.int3_params[AC_global_grid_n]) +//#define globalMeshM // Placeholder +//#define localMeshN // Placeholder +//#define localMeshM // Placeholder +//#define localMeshN_min // Placeholder +//#define globalMeshN_min // Placeholder #define d_multigpu_offset (d_mesh_info.int3_params[AC_multigpu_offset]) +//#define d_multinode_offset (d_mesh_info.int3_params[AC_multinode_offset]) // Placeholder #include "kernels/boundconds.cuh" #include "kernels/integration.cuh" #include "kernels/reductions.cuh"