Add span_alloc for the size of gpu data allocation

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++ );

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.

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<<i;
         gpu_task->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;

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");
         }
 

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

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

tests/runtime/cuda/CMakeLists.txt

@@ -31,6 +31,10 @@ 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")
 
+  parsec_addtest_executable(C alloc_size SOURCES alloc_size_main.c)
+  target_include_directories(alloc_size PRIVATE $<$<NOT:${PARSEC_BUILD_INPLACE}>:${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)

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()

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 <stdint.h>
+#if defined(DISTRIBUTED)
+#include <mpi.h>
+#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;
+}
+
+%}