diff --git a/tests/runtime/cuda/CMakeLists.txt b/tests/runtime/cuda/CMakeLists.txt index 6d1479fb1..c464fd791 100644 --- a/tests/runtime/cuda/CMakeLists.txt +++ b/tests/runtime/cuda/CMakeLists.txt @@ -31,6 +31,11 @@ if(PARSEC_HAVE_CUDA) target_include_directories(testing_get_best_device PRIVATE $<$:${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 $<$:${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) diff --git a/tests/runtime/cuda/reallocate_tile_example.jdf b/tests/runtime/cuda/reallocate_tile_example.jdf new file mode 100644 index 000000000..1355772ed --- /dev/null +++ b/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 +#include +#include + +#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 diff --git a/tests/runtime/cuda/reallocate_tile_example_main.c b/tests/runtime/cuda/reallocate_tile_example_main.c new file mode 100644 index 000000000..93c5f39a4 --- /dev/null +++ b/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 +#endif + +#include +#include +#include + +/* 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; +}