Add example for memory reallocation on GPU

tests/runtime/cuda/CMakeLists.txt

@@ -31,6 +31,11 @@ if(PARSEC_HAVE_CUDA)
   target_include_directories(testing_get_best_device PRIVATE $<$<NOT:${PARSEC_BUILD_INPLACE}>:${CMAKE_CURRENT_SOURCE_DIR}>)
   target_ptg_sources(testing_get_best_device PRIVATE "get_best_device_check.jdf")
 
+  # Testing for GPU tile reallocation path used by HICMA runtime decision
+  parsec_addtest_executable(C reallocate_tile_example SOURCES "reallocate_tile_example_main.c")
+  target_include_directories(reallocate_tile_example PRIVATE $<$<NOT:${PARSEC_BUILD_INPLACE}>:${CMAKE_CURRENT_SOURCE_DIR}>)
+  target_ptg_sources(reallocate_tile_example PRIVATE "reallocate_tile_example.jdf")
+
   if(CMAKE_CUDA_COMPILER)
     set_source_files_properties(ping_kernel.cu PROPERTIES LANGUAGE CUDA)
     list(APPEND PINGPONG_GPU_KERNEL_SOURCES ping_kernel.cu)

tests/runtime/cuda/reallocate_tile_example.jdf

@@ -0,0 +1,213 @@
+extern "C" %{
+#include "parsec/data_distribution.h"
+#include "parsec/data_dist/matrix/matrix.h"
+#include "parsec/data_internal.h"
+#include "parsec/mca/device/cuda/device_cuda_internal.h"
+#include "parsec/utils/zone_malloc.h"
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+
+#define ITER_COUNT 10
+
+static int hicma_reallocate_tile_on_gpu(parsec_device_cuda_module_t *cuda_device,
+                                        parsec_data_copy_t *active_copy,
+                                        void *expected_ptr,
+                                        size_t target_bytes,
+                                        void **new_ptr_out)
+{
+    void *old_ptr = NULL;
+    void *new_ptr = NULL;
+
+    if( NULL == active_copy ) {
+        return -1;
+    }
+    if( NULL != expected_ptr && expected_ptr != active_copy->device_private ) {
+        return -1;
+    }
+
+    old_ptr = active_copy->device_private;
+    if( NULL != old_ptr ) {
+        zone_free(cuda_device->super.memory, old_ptr);
+        active_copy->device_private = NULL;
+    }
+
+    new_ptr = zone_malloc(cuda_device->super.memory, target_bytes);
+    if( NULL == new_ptr ) {
+        return -1;
+    }
+
+    active_copy->device_private = new_ptr;
+    *new_ptr_out = new_ptr;
+    return 0;
+}
+
+static int hicma_reallocate_tile_on_cpu(parsec_data_copy_t *cpu_copy,
+                                        size_t target_bytes)
+{
+    void *new_ptr = NULL;
+
+    if( NULL == cpu_copy ) {
+        return 0;
+    }
+
+    new_ptr = malloc(target_bytes);
+    if( NULL == new_ptr ) {
+        return -1;
+    }
+
+    if( NULL != cpu_copy->device_private ) {
+        free(cpu_copy->device_private);
+    }
+    cpu_copy->device_private = new_ptr;
+    return 0;
+}
+%}
+
+%option no_taskpool_instance = false
+
+descA [type = "parsec_tiled_matrix_t *"]
+OUT   [type = "int *"]
+
+TASK(k)
+k = 0 .. 10
+
+: descA(0, 0)
+
+RW A <- (k == 0) ? descA(0, 0) : A TASK(k-1)
+     -> (k < 10) ? A TASK(k+1)
+     -> (k == 10) ? descA(0, 0)
+CTL T -> (k == 10) ? T TEST(0)
+
+BODY [type=CUDA]
+{
+    parsec_data_copy_t *c_copy = this_task->data._f_A.data_out;
+    parsec_data_copy_t *c_dev_copy = NULL;
+    parsec_data_copy_t *active_copy = NULL;
+    parsec_data_copy_t *cpu_copy = NULL;
+    void *new_A = NULL;
+    void *A_ptr = NULL;
+    double *host_init = NULL;
+    size_t target_bytes = 0;
+    const size_t base_count = (size_t)descA->mb * (size_t)descA->nb;
+    const size_t grown_count = base_count * 2U;
+
+    if( NULL == c_copy || NULL == c_copy->original ) {
+        *OUT = *OUT + 1;
+        fprintf(stderr, "TASK(%d): missing data copy\n", k);
+        return PARSEC_HOOK_RETURN_DONE;
+    }
+
+    c_dev_copy = PARSEC_DATA_GET_COPY(c_copy->original, c_copy->device_index);
+    cpu_copy = PARSEC_DATA_GET_COPY(c_copy->original, 0);
+    active_copy = (NULL != c_dev_copy) ? c_dev_copy : c_copy;
+    A_ptr = c_copy->device_private;
+
+    if( k == 0 ) {
+        void *new_cpu_stage_ptr = NULL;
+
+        target_bytes = grown_count * sizeof(double);
+        if( 0 != hicma_reallocate_tile_on_gpu(cuda_device, active_copy, A_ptr, target_bytes, &new_A) ) {
+            *OUT = *OUT + 1;
+            fprintf(stderr, "TASK(%d): GPU reallocation to %zu doubles failed\n", k, grown_count);
+            return PARSEC_HOOK_RETURN_DONE;
+        }
+        if( c_copy != active_copy ) {
+            c_copy->device_private = new_A;
+        }
+        /*
+         * In this PaRSEC version, data transfer size follows original->span.
+         * Update span after growth and provide a sufficiently large CPU staging
+         * buffer to keep subsequent D2H copies valid.
+         */
+        c_copy->original->span = target_bytes;
+        if( NULL != cpu_copy ) {
+            new_cpu_stage_ptr = malloc(target_bytes);
+            if( NULL == new_cpu_stage_ptr ) {
+                *OUT = *OUT + 1;
+                fprintf(stderr, "TASK(%d): CPU stage buffer reallocation failed\n", k);
+                return PARSEC_HOOK_RETURN_DONE;
+            }
+            cpu_copy->device_private = new_cpu_stage_ptr;
+        }
+        host_init = (double *)malloc(target_bytes);
+        if( NULL == host_init ) {
+            *OUT = *OUT + 1;
+            fprintf(stderr, "TASK(%d): host initialization buffer allocation failed\n", k);
+            return PARSEC_HOOK_RETURN_DONE;
+        }
+        for(size_t i = 0; i < grown_count; i++) {
+            host_init[i] = 1000.0;
+        }
+        if( cudaSuccess != cudaMemcpyAsync(new_A, host_init, target_bytes,
+                                           cudaMemcpyHostToDevice, parsec_body.stream) ) {
+            *OUT = *OUT + 1;
+            fprintf(stderr, "TASK(%d): set GPU values to 1000 failed\n", k);
+            free(host_init);
+            return PARSEC_HOOK_RETURN_DONE;
+        }
+        free(host_init);
+    }
+
+    if( k == ITER_COUNT ) {
+        parsec_data_copy_t mock_cpu_copy;
+
+        target_bytes = grown_count * sizeof(double);
+        memset(&mock_cpu_copy, 0, sizeof(mock_cpu_copy));
+        mock_cpu_copy.device_private = malloc(base_count * sizeof(double));
+        if( NULL == mock_cpu_copy.device_private ) {
+            *OUT = *OUT + 1;
+            fprintf(stderr, "TASK(%d): initial mock CPU allocation failed\n", k);
+            return PARSEC_HOOK_RETURN_DONE;
+        }
+
+        if( 0 != hicma_reallocate_tile_on_cpu(&mock_cpu_copy, target_bytes) ) {
+            *OUT = *OUT + 1;
+            fprintf(stderr, "TASK(%d): CPU reallocation to %zu doubles failed\n", k, grown_count);
+            free(mock_cpu_copy.device_private);
+            return PARSEC_HOOK_RETURN_DONE;
+        }
+        memset(mock_cpu_copy.device_private, 1, target_bytes);
+        free(mock_cpu_copy.device_private);
+    }
+}
+END
+
+BODY
+{
+    *OUT = *OUT + 1;
+    fprintf(stderr, "TASK(%d) ran on CPU; expected CUDA BODY\n", k);
+}
+END
+
+TEST(k)
+k = 0 .. 0
+
+: descA(0, 0)
+
+READ A <- descA(0, 0)
+CTL T  <- T TASK(10)
+
+BODY
+{
+    size_t base_count = (size_t)descA->mb * (size_t)descA->nb;
+    size_t grown_count = base_count * 2U;
+    double *host_check = (double *)A;
+
+    if( NULL == host_check ) {
+        *OUT = *OUT + 1;
+        fprintf(stderr, "TEST(%d): missing CPU pointer\n", k);
+        return PARSEC_HOOK_RETURN_DONE;
+    }
+
+    for(size_t i = 0; i < grown_count; i++) {
+        if( host_check[i] != 1000.0 ) {
+            *OUT = *OUT + 1;
+            fprintf(stderr, "TEST(%d): value mismatch at index %zu (got %.17g)\n",
+                    k, i, host_check[i]);
+            return PARSEC_HOOK_RETURN_DONE;
+        }
+    }
+}
+END

tests/runtime/cuda/reallocate_tile_example_main.c

@@ -0,0 +1,89 @@
+#include "parsec.h"
+#include "parsec/data_distribution.h"
+#include "parsec/data_dist/matrix/matrix.h"
+#include "parsec/data_dist/matrix/two_dim_rectangle_cyclic.h"
+
+#include "reallocate_tile_example.h"
+
+#if defined(DISTRIBUTED)
+#include <mpi.h>
+#endif
+
+#include <assert.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+/* The generated internal taskpool class derives from this public taskpool type. */
+PARSEC_OBJ_CLASS_INSTANCE(parsec_reallocate_tile_example_taskpool_t, parsec_taskpool_t,
+                          NULL, NULL);
+
+int main(int argc, char **argv)
+{
+    parsec_context_t *parsec = NULL;
+    parsec_reallocate_tile_example_taskpool_t *tp = NULL;
+    parsec_matrix_block_cyclic_t dcA;
+    int rank = 0;
+    int ret = 0;
+    const int tile_size = 32; /* 32x32 = 1024 doubles in one tile */
+
+#if defined(DISTRIBUTED)
+    {
+        int provided;
+        MPI_Init_thread(NULL, NULL, MPI_THREAD_SERIALIZED, &provided);
+    }
+    MPI_Comm_rank(MPI_COMM_WORLD, &rank);
+#endif
+
+    parsec = parsec_init(1, &argc, &argv);
+    if( NULL == parsec ) {
+        return EXIT_FAILURE;
+    }
+
+    /* This test validates a CUDA BODY path and should be skipped without GPU. */
+    if( 0 == parsec_context_query(parsec, PARSEC_CONTEXT_QUERY_DEVICES, PARSEC_DEV_CUDA) ) {
+        parsec_warning("reallocate_tile_example requires at least one CUDA device");
+        printf("TEST SKIPPED\n");
+        parsec_fini(&parsec);
+#if defined(DISTRIBUTED)
+        MPI_Finalize();
+#endif
+        return EXIT_SUCCESS;
+    }
+
+    parsec_matrix_block_cyclic_init(&dcA, PARSEC_MATRIX_DOUBLE, PARSEC_MATRIX_TILE,
+                                    rank,
+                                    tile_size, tile_size,
+                                    tile_size, tile_size,
+                                    0, 0,
+                                    tile_size, tile_size,
+                                    1, 1, 1, 1, 0, 0);
+    dcA.mat = parsec_data_allocate((size_t)dcA.super.nb_local_tiles *
+                                   (size_t)dcA.super.bsiz *
+                                   (size_t)parsec_datadist_getsizeoftype(dcA.super.mtype));
+    assert(NULL != dcA.mat);
+    parsec_data_collection_set_key((parsec_data_collection_t *)&dcA, "dcA_reallocate");
+
+    tp = parsec_reallocate_tile_example_new((parsec_tiled_matrix_t *)&dcA, &ret);
+    assert(NULL != tp);
+
+    parsec_context_add_taskpool(parsec, &tp->super);
+    parsec_context_start(parsec);
+    parsec_context_wait(parsec);
+    parsec_taskpool_free(&tp->super);
+
+    parsec_data_free(dcA.mat);
+    parsec_tiled_matrix_destroy((parsec_tiled_matrix_t *)&dcA);
+    parsec_fini(&parsec);
+
+#if defined(DISTRIBUTED)
+    MPI_Finalize();
+#endif
+
+    if( 0 == ret ) {
+        printf("reallocate_tile_example: TEST PASSED\n");
+        return EXIT_SUCCESS;
+    }
+
+    printf("reallocate_tile_example: TEST FAILED (%d)\n", ret);
+    return EXIT_FAILURE;
+}