diff --git a/parsec/data.c b/parsec/data.c index 29770d593..c0b90db8d 100644 --- a/parsec/data.c +++ b/parsec/data.c @@ -75,6 +75,7 @@ static void parsec_data_construct(parsec_data_t* obj ) obj->preferred_device = -1; obj->key = 0; obj->span = 0; + obj->span_alloc = 0; obj->nb_copies = 0; for( uint32_t i = 0; i < parsec_nb_devices; obj->device_copies[i] = NULL, i++ ); diff --git a/parsec/data_internal.h b/parsec/data_internal.h index 30662c753..4f7a36502 100644 --- a/parsec/data_internal.h +++ b/parsec/data_internal.h @@ -39,7 +39,10 @@ struct parsec_data_s { int32_t nb_copies; /* How many valid copies are attached to this data */ parsec_data_key_t key; struct parsec_data_collection_s* dc; - size_t span; /* size in bytes of the memory layout */ + size_t span; /* logical size in bytes of the memory layout */ + size_t span_alloc; /* allocated capacity in bytes for device buffers. + * 0 means "unset" and should fallback to span. + */ struct parsec_data_copy_s *device_copies[]; /* this array allocated according to the number of devices * (parsec_nb_devices). It points to the most recent * version of the data. diff --git a/parsec/interfaces/dtd/insert_function.c b/parsec/interfaces/dtd/insert_function.c index 2fecfa1a0..1dbf36ba8 100644 --- a/parsec/interfaces/dtd/insert_function.c +++ b/parsec/interfaces/dtd/insert_function.c @@ -2409,6 +2409,9 @@ static parsec_hook_return_t parsec_dtd_gpu_task_submit(parsec_execution_stream_t gpu_task->pushout |= 1<flow_info[i].flow = dtd_tc->super.in[i]; gpu_task->flow_info[i].flow_span = this_task->data[i].data_in->original->span; + gpu_task->flow_info[i].flow_span_alloc = (0 == this_task->data[i].data_in->original->span_alloc) ? + gpu_task->flow_info[i].flow_span : + this_task->data[i].data_in->original->span_alloc; } parsec_device_module_t *device = this_task->selected_device; diff --git a/parsec/interfaces/ptg/ptg-compiler/jdf2c.c b/parsec/interfaces/ptg/ptg-compiler/jdf2c.c index f63eb0e8e..39b4ff623 100644 --- a/parsec/interfaces/ptg/ptg-compiler/jdf2c.c +++ b/parsec/interfaces/ptg/ptg-compiler/jdf2c.c @@ -6619,6 +6619,7 @@ static void jdf_generate_code_hook_gpu(const jdf_t *jdf, jdf_def_list_t *stage_in_property; jdf_def_list_t *stage_out_property; jdf_def_list_t *size_property; + jdf_def_list_t *alloc_size_property; jdf_def_list_t *desc_property; jdf_def_list_t *device_property; jdf_def_list_t *prop; @@ -6877,14 +6878,17 @@ static void jdf_generate_code_hook_gpu(const jdf_t *jdf, sprintf(sa->ptr, "%s.size", fl->varname); jdf_find_property(body->properties, sa->ptr, &size_property); + sprintf(sa->ptr, "%s.alloc_size", fl->varname); + jdf_find_property(body->properties, sa->ptr, &alloc_size_property); if(fl->flow_flags & JDF_FLOW_TYPE_CTL) { - if(size_property != NULL) { - fprintf(stderr, "Error: specifying GPU buffer size for CTL flow %s at line %d\n", + if(size_property != NULL || alloc_size_property != NULL) { + fprintf(stderr, "Error: specifying GPU size/alloc_size for CTL flow %s at line %d\n", fl->varname, JDF_OBJECT_LINENO(fl)); exit(-1); } coutput(" gpu_task->flow_info[%d].flow_span = 0;\n", di); + coutput(" gpu_task->flow_info[%d].flow_span_alloc = 0;\n", di); } else { coutput(" // A shortcut to check if the flow exists\n"); coutput(" if (gpu_task->ec->data[%d].data_in != NULL) {\n", di); @@ -6899,6 +6903,13 @@ static void jdf_generate_code_hook_gpu(const jdf_t *jdf, } } + if(alloc_size_property == NULL){ + coutput(" gpu_task->flow_info[%d].flow_span_alloc = (0 == gpu_task->ec->data[%d].data_in->original->span_alloc) ? gpu_task->flow_info[%d].flow_span : gpu_task->ec->data[%d].data_in->original->span_alloc;\n", di, di, di, di); + } else { + coutput(" gpu_task->flow_info[%d].flow_span_alloc = %s;\n", + di, dump_expr((void**)alloc_size_property->expr, &info)); + } + coutput(" assert(gpu_task->flow_info[%d].flow_span <= gpu_task->flow_info[%d].flow_span_alloc);\n", di, di); coutput("}\n"); } diff --git a/parsec/mca/device/device_gpu.c b/parsec/mca/device/device_gpu.c index 67e7ca6c3..6af8230b5 100644 --- a/parsec/mca/device/device_gpu.c +++ b/parsec/mca/device/device_gpu.c @@ -748,6 +748,9 @@ parsec_device_data_advise(parsec_device_module_t *dev, parsec_data_t *data, int gpu_task->nb_flows = 1; gpu_task->flow_info[0].flow = &parsec_device_data_prefetch_flow; gpu_task->flow_info[0].flow_span = data->device_copies[ data->owner_device ]->original->span; + gpu_task->flow_info[0].flow_span_alloc = (0 == data->device_copies[ data->owner_device ]->original->span_alloc) ? + gpu_task->flow_info[0].flow_span : + data->device_copies[ data->owner_device ]->original->span_alloc; gpu_task->stage_in = parsec_default_gpu_stage_in; gpu_task->stage_out = parsec_default_gpu_stage_out; PARSEC_DEBUG_VERBOSE(20, parsec_debug_output, "Retain data copy %p [ref_count %d]", @@ -1244,6 +1247,7 @@ parsec_device_data_reserve_space( parsec_device_gpu_module_t* gpu_device, /* Skip CTL flows only */ if(PARSEC_FLOW_ACCESS_NONE == (PARSEC_FLOW_ACCESS_MASK & flow->flow_flags)) { gpu_task->flow_info[i].flow_span = 0; /* assume there is nothing to transfer to the GPU */ + gpu_task->flow_info[i].flow_span_alloc = 0; continue; } @@ -1315,15 +1319,21 @@ parsec_device_data_reserve_space( parsec_device_gpu_module_t* gpu_device, #if !defined(PARSEC_GPU_ALLOC_PER_TILE) gpu_elem = PARSEC_OBJ_NEW(parsec_data_copy_t); + size_t master_span_alloc = (0 == master->span_alloc) ? master->span : master->span_alloc; + size_t flow_span_alloc = (0 == gpu_task->flow_info[i].flow_span_alloc) ? gpu_task->flow_info[i].flow_span : gpu_task->flow_info[i].flow_span_alloc; + size_t alloc_span = (flow_span_alloc > master_span_alloc) ? flow_span_alloc : master_span_alloc; + assert(master->span <= master_span_alloc); + assert(gpu_task->flow_info[i].flow_span <= flow_span_alloc); PARSEC_DEBUG_VERBOSE(20, parsec_gpu_output_stream, - "GPU[%d:%s]:%s: Allocate GPU copy %p sz %zu [ref_count %d] for data %p", + "GPU[%d:%s]:%s: Allocate GPU copy %p flow_sz %zu alloc_sz %zu [ref_count %d] for data %p", gpu_device->super.device_index, gpu_device->super.name, task_name, - gpu_elem, gpu_task->flow_info[i].flow_span, gpu_elem->super.super.obj_reference_count, master); + gpu_elem, gpu_task->flow_info[i].flow_span, alloc_span, + gpu_elem->super.super.obj_reference_count, master); gpu_elem->flags = PARSEC_DATA_FLAG_PARSEC_OWNED | PARSEC_DATA_FLAG_PARSEC_MANAGED; malloc_data: copy_readers_update = 0; assert(0 != (gpu_elem->flags & PARSEC_DATA_FLAG_PARSEC_OWNED) ); - gpu_elem->device_private = zone_malloc(gpu_device->memory, gpu_task->flow_info[i].flow_span); + gpu_elem->device_private = zone_malloc(gpu_device->memory, alloc_span); gpu_elem->arena_chunk = (parsec_arena_chunk_t *)gpu_device->memory; if( NULL == gpu_elem->device_private ) { #endif @@ -1537,7 +1547,7 @@ parsec_device_data_reserve_space( parsec_device_gpu_module_t* gpu_device, parsec_profiling_trace_flags(gpu_device->exec_stream[0]->profiling, parsec_gpu_allocate_memory_key, (int64_t)gpu_elem->device_private, gpu_device->super.device_index, - &gpu_task->flow_info[i].flow_span, PARSEC_PROFILING_EVENT_COUNTER|PARSEC_PROFILING_EVENT_HAS_INFO); + &alloc_span, PARSEC_PROFILING_EVENT_COUNTER|PARSEC_PROFILING_EVENT_HAS_INFO); } #endif #else @@ -2335,6 +2345,9 @@ parsec_device_send_transfercomplete_cmd_to_device(parsec_data_copy_t *copy, gpu_task->nb_flows = 1; gpu_task->flow_info[0].flow = &parsec_device_d2d_complete_flow; gpu_task->flow_info[0].flow_span = copy->original->span; + gpu_task->flow_info[0].flow_span_alloc = (0 == copy->original->span_alloc) ? + gpu_task->flow_info[0].flow_span : + copy->original->span_alloc; gpu_task->stage_in = parsec_default_gpu_stage_in; gpu_task->stage_out = parsec_default_gpu_stage_out; gpu_task->ec->data[0].data_in = copy; /* We need to set not-null in data_in, so that the fake flow is diff --git a/parsec/mca/device/device_gpu.h b/parsec/mca/device/device_gpu.h index b36a40718..f245df6ac 100644 --- a/parsec/mca/device/device_gpu.h +++ b/parsec/mca/device/device_gpu.h @@ -101,6 +101,10 @@ typedef struct parsec_gpu_flow_info_s { * size of the data, for all the other copies this should be the amount of memory * needed on the device. */ + size_t flow_span_alloc; /* allocation span on the device for this flow. + * Defaults to original->span_alloc when not overridden; + * if that is 0 it falls back to flow_span. + */ parsec_data_collection_t *flow_dc; /* the data collection from which the data originates. When the data copy is local, the data * collection can be accessed via the data_t, but for all copies coming from the network there * is no known data collection. Thus, for such cases the DSL need to provide a reference to the diff --git a/tests/runtime/cuda/CMakeLists.txt b/tests/runtime/cuda/CMakeLists.txt index 6d1479fb1..11c0c4030 100644 --- a/tests/runtime/cuda/CMakeLists.txt +++ b/tests/runtime/cuda/CMakeLists.txt @@ -31,6 +31,10 @@ 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") + parsec_addtest_executable(C alloc_size SOURCES alloc_size_main.c) + target_include_directories(alloc_size PRIVATE $<$:${CMAKE_CURRENT_SOURCE_DIR}>) + target_ptg_sources(alloc_size PRIVATE "alloc_size.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/Testings.cmake b/tests/runtime/cuda/Testings.cmake index 03c85ed36..ff140cca8 100644 --- a/tests/runtime/cuda/Testings.cmake +++ b/tests/runtime/cuda/Testings.cmake @@ -9,6 +9,9 @@ if(PARSEC_HAVE_CUDA) if(TARGET stage) parsec_addtest_cmd(runtime/cuda/stage:gpu ${SHM_TEST_CMD_LIST} ${CTEST_CUDA_LAUNCHER_OPTIONS} runtime/cuda/stage --mca device_cuda_enabled 2 --mca device_show_statistics 1) endif() + if(TARGET alloc_size) + parsec_addtest_cmd(runtime/cuda/alloc_size:gpu ${SHM_TEST_CMD_LIST} ${CTEST_CUDA_LAUNCHER_OPTIONS} runtime/cuda/alloc_size --mca device_cuda_enabled 1 --mca device_show_statistics 1) + endif() if(TARGET ptg_pingpong) parsec_addtest_cmd(runtime/cuda/ptg_pingpong:gpu ${SHM_TEST_CMD_LIST} ${CTEST_CUDA_LAUNCHER_OPTIONS} runtime/cuda/ptg_pingpong --mca device_cuda_enabled 2 --mca device_show_statistics 1) endif() diff --git a/tests/runtime/cuda/alloc_size.jdf b/tests/runtime/cuda/alloc_size.jdf new file mode 100644 index 000000000..8c8703362 --- /dev/null +++ b/tests/runtime/cuda/alloc_size.jdf @@ -0,0 +1,177 @@ +extern "C" %{ +/* + * Copyright (c) 2026 The University of Tennessee and The University + * of Tennessee Research Foundation. All rights + * reserved. + */ +#include "parsec/parsec_config.h" + +#include "parsec/data_distribution.h" +#include "parsec/data_dist/matrix/matrix.h" +#include "parsec/data_dist/matrix/two_dim_rectangle_cyclic.h" +#include +#if defined(DISTRIBUTED) +#include +#endif + +#include "alloc_size.h" + +static inline void alloc_size_expected_spans(int m, int n, size_t total, + size_t *logical, size_t *span_alloc_set) +{ + /* + * Per-tile coverage for allocation sizing paths: + * case 0 ((m+n)%3==0): explicit over-allocation request (span_alloc > span). + * Verifies non-zero span_alloc is preserved end-to-end. + * case 1 ((m+n)%3==1): partial logical span with span_alloc == 0. + * Verifies fallback path where allocator must use the logical span. + * case 2 ((m+n)%3==2): full logical span with span_alloc == 0. + * Verifies zero span_alloc remains valid for full-span transfers too. + */ + const int tile_case = (m + n) % 3; + if(tile_case == 0) { + *logical = total / 2; + *span_alloc_set = *logical + 64; + if(*span_alloc_set > total) *span_alloc_set = total; + } else if(tile_case == 1) { + *logical = (3 * total) / 4; + *span_alloc_set = 0; + } else { + *logical = total; + *span_alloc_set = 0; + } +} + +%} + +%option no_taskpool_instance = true + +descA [type = "parsec_tiled_matrix_t*"] +info [type = "int*"] + +set_sizes_cpu(m, n) + +m = 0 .. descA->mt-1 +n = 0 .. descA->nt-1 + +: descA(m, n) + +RW A <- descA(m, n) + -> A task_gpu(m, n) + +BODY +{ + size_t total = (size_t)descA->mb * (size_t)descA->nb * parsec_datadist_getsizeoftype(descA->mtype); + size_t logical = 0; + size_t span_alloc_set = 0; + uint32_t marker = (uint32_t)(0x11110000u | ((uint32_t)(m & 0xFF) << 8) | (uint32_t)(n & 0xFF)); + int nelems = descA->mb * descA->nb; + + alloc_size_expected_spans(m, n, total, &logical, &span_alloc_set); + if(nelems > 0) { + ((uint32_t*)A)[0] = marker; + ((uint32_t*)A)[nelems - 1] = marker ^ 0x00FF00FFu; + } + + this_task->data._f_A.data_out->original->span = logical; + this_task->data._f_A.data_out->original->span_alloc = span_alloc_set; +} +END + +task_gpu(m, n) + +m = 0 .. descA->mt-1 +n = 0 .. descA->nt-1 + +: descA(m, n) + +RW A <- A set_sizes_cpu(m, n) + -> A check(m, n) + +BODY [type=CUDA] +{ + (void)A; +} +END + +check(m, n) + +m = 0 .. descA->mt-1 +n = 0 .. descA->nt-1 + +: descA(m, n) + +READ A <- A task_gpu(m, n) + +BODY +{ + size_t total = (size_t)descA->mb * (size_t)descA->nb * parsec_datadist_getsizeoftype(descA->mtype); + size_t logical = 0; + size_t span_alloc_set = 0; + uint32_t marker = (uint32_t)(0x11110000u | ((uint32_t)(m & 0xFF) << 8) | (uint32_t)(n & 0xFF)); + int nelems = descA->mb * descA->nb; + + alloc_size_expected_spans(m, n, total, &logical, &span_alloc_set); + + if(this_task->data._f_A.data_in->original->span != logical) { + info[es->th_id]++; + } + if(this_task->data._f_A.data_in->original->span_alloc != span_alloc_set) { + info[es->th_id]++; + } + if(nelems > 0) { + if(((uint32_t*)A)[0] != marker) { + info[es->th_id]++; + } + if(((uint32_t*)A)[nelems - 1] != (marker ^ 0x00FF00FFu)) { + info[es->th_id]++; + } + } +} +END + +extern "C" %{ + +static parsec_taskpool_t* parsec_alloc_size_test_New(parsec_tiled_matrix_t *dcA, int *info) +{ + parsec_alloc_size_taskpool_t *taskpool = parsec_alloc_size_new(dcA, info); + parsec_matrix_arena_datatype_define_type(&taskpool->arenas_datatypes[PARSEC_alloc_size_DEFAULT_ADT_IDX], + parsec_datatype_int32_t, PARSEC_MATRIX_FULL, + 1, dcA->mb, dcA->nb, dcA->mb, + PARSEC_ARENA_ALIGNMENT_SSE, -1); + return (parsec_taskpool_t*)taskpool; +} + +static void __parsec_taskpool_alloc_size_destructor(parsec_alloc_size_taskpool_t *tp) +{ + parsec_matrix_arena_datatype_destruct_free_type(&tp->arenas_datatypes[PARSEC_alloc_size_DEFAULT_ADT_IDX]); +} + +PARSEC_OBJ_CLASS_INSTANCE(parsec_alloc_size_taskpool_t, parsec_taskpool_t, + NULL, __parsec_taskpool_alloc_size_destructor); + +int parsec_alloc_size_test(parsec_context_t *parsec, parsec_tiled_matrix_t *A) +{ + int info = 0; + int nb_threads = parsec->virtual_processes[0]->nb_cores; + int *info_tmp = (int *)calloc((size_t)nb_threads, sizeof(int)); + parsec_taskpool_t *tp = parsec_alloc_size_test_New(A, info_tmp); + + parsec_enqueue(parsec, tp); + parsec_context_start(parsec); + parsec_context_wait(parsec); + parsec_taskpool_free(tp); + + for(int i = 1; i < nb_threads; i++) info_tmp[0] += info_tmp[i]; + +#if defined(DISTRIBUTED) + MPI_Allreduce(&info_tmp[0], &info, 1, MPI_INT, MPI_SUM, MPI_COMM_WORLD); +#else + info = info_tmp[0]; +#endif + + free(info_tmp); + return info; +} + +%} diff --git a/tests/runtime/cuda/alloc_size_main.c b/tests/runtime/cuda/alloc_size_main.c new file mode 100644 index 000000000..1fa873dbc --- /dev/null +++ b/tests/runtime/cuda/alloc_size_main.c @@ -0,0 +1,103 @@ +/* + * Copyright (c) 2026 The University of Tennessee and The University + * of Tennessee Research Foundation. All rights + * reserved. + */ +#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 "parsec/data_internal.h" + +#if defined(DISTRIBUTED) +#include +#endif + +int parsec_alloc_size_test(parsec_context_t *parsec, parsec_tiled_matrix_t *A); + +int main(int argc, char *argv[]) +{ + parsec_context_t *parsec = NULL; + parsec_matrix_block_cyclic_t dcA; + int rank = 0, size = 1; + int info = 0; + +#if defined(DISTRIBUTED) + { + int provided; + MPI_Init_thread(NULL, NULL, MPI_THREAD_SERIALIZED, &provided); + } + MPI_Comm_size(MPI_COMM_WORLD, &size); + MPI_Comm_rank(MPI_COMM_WORLD, &rank); +#endif + + parsec = parsec_init(-1, &argc, &argv); + if(NULL == parsec) return EXIT_FAILURE; + + if(size != 1) { + if(rank == 0) fprintf(stderr, "alloc_size test supports single-node runtime launch only\n"); + parsec_fini(&parsec); +#if defined(DISTRIBUTED) + MPI_Finalize(); +#endif + return EXIT_FAILURE; + } + + if(parsec_context_query(parsec, PARSEC_CONTEXT_QUERY_DEVICES, PARSEC_DEV_CUDA) == 0) { + if(rank == 0) { + parsec_warning("This test requires at least one CUDA device"); + printf("TEST SKIPPED\n"); + } + parsec_fini(&parsec); +#if defined(DISTRIBUTED) + MPI_Finalize(); +#endif + return -PARSEC_ERR_DEVICE; + } + + parsec_matrix_block_cyclic_init(&dcA, PARSEC_MATRIX_INTEGER, PARSEC_MATRIX_TILE, + rank, 32, 32, 128, 128, 0, 0, + 128, 128, 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)); + parsec_data_collection_set_key((parsec_data_collection_t*)&dcA, "dcA_alloc_size"); + + { + int *vals = (int*)dcA.mat; + size_t count = (size_t)dcA.super.nb_local_tiles * (size_t)dcA.super.bsiz; + for(size_t i = 0; i < count; i++) vals[i] = 0x11111111; + } + + { + parsec_data_collection_t *dc = (parsec_data_collection_t*)&dcA; + for(int m = 0; m < dcA.super.mt; m++) { + for(int n = 0; n < dcA.super.nt; n++) { + parsec_data_t *data = dc->data_of(dc, m, n); + if(NULL == data) continue; + if(data->span_alloc != 0) { + fprintf(stderr, "default span_alloc is expected 0 but got %zu at (%d,%d)\n", + data->span_alloc, m, n); + info++; + } + } + } + } + + if(info == 0) { + info = parsec_alloc_size_test(parsec, (parsec_tiled_matrix_t *)&dcA); + } + + if(rank == 0) { + if(info == 0) printf("TEST PASSED\n"); + else printf("TEST FAILED (%d errors)\n", info); + } + + parsec_data_free(dcA.mat); + parsec_tiled_matrix_destroy((parsec_tiled_matrix_t*)&dcA); + parsec_fini(&parsec); +#if defined(DISTRIBUTED) + MPI_Finalize(); +#endif + return (0 == info) ? EXIT_SUCCESS : EXIT_FAILURE; +}