diff --git a/parsec/mca/device/cuda/device_cuda_module.c b/parsec/mca/device/cuda/device_cuda_module.c index 932d968e8..e9ba8ae95 100644 --- a/parsec/mca/device/cuda/device_cuda_module.c +++ b/parsec/mca/device/cuda/device_cuda_module.c @@ -343,6 +343,58 @@ static int parsec_cuda_memcpy_async(struct parsec_device_gpu_module_s *gpu, stru return PARSEC_SUCCESS; } +#if CUDART_VERSION >= 12080 +/** + * @brief Native implementation of parsec_device_memcpy_multi_async_fn_t using + * cudaMemcpyBatchAsync (introduced in CUDA 12.8). @p dsts/@p srcs/@p sizes are + * forwarded to cudaMemcpyBatchAsync as-is -- no repacking -- since its parameter + * shape is exactly the parallel-array layout parsec_device_memcpy_multi_async_fn_t + * already uses. @p directions is unused here: cudaMemcpyBatchAsync infers each + * copy's kind from the pointers themselves (unified addressing), unlike the + * single-item cudaMemcpyAsync which needs an explicit cudaMemcpyKind. + * All items share the same cudaMemcpyAttributes (stream-ordered access, matching + * the semantics of the individual cudaMemcpyAsync calls this replaces); location + * hints are left unset since they only matter for managed memory, which PaRSEC + * does not use for these buffers. The cudaMemcpyBatchAsync signature dropped its + * failIdx parameter in CUDA 13.0; guard both forms. + */ +static int parsec_cuda_memcpy_multi_async(struct parsec_device_gpu_module_s *gpu, struct parsec_gpu_exec_stream_s *gpu_stream, + void **dsts, void **srcs, size_t *sizes, + parsec_device_transfer_direction_t *directions, int nb_items) +{ + parsec_cuda_exec_stream_t *cuda_stream = (parsec_cuda_exec_stream_t *)gpu_stream; + cudaMemcpyAttributes attrs = {0}; + size_t attrsIdx = 0; + cudaError_t cudaStatus; + + (void)gpu; + + if( 0 == nb_items ) + return PARSEC_SUCCESS; + + if (1 == nb_items) { + // short-cut to the single-item memcpy_async implementation, which takes + // into account the direction parameter and is more efficient than cudaMemcpyBatchAsync for a single item. + return parsec_cuda_memcpy_async(gpu, gpu_stream, dsts[0], srcs[0], sizes[0], directions[0]); + } + + attrs.srcAccessOrder = cudaMemcpySrcAccessOrderStream; + +#if CUDART_VERSION >= 13000 + cudaStatus = cudaMemcpyBatchAsync(dsts, (const void *const *)srcs, sizes, (size_t)nb_items, + &attrs, &attrsIdx, 1 /* numAttrs */, + cuda_stream->cuda_stream); +#else + cudaStatus = cudaMemcpyBatchAsync(dsts, srcs, sizes, (size_t)nb_items, + &attrs, &attrsIdx, 1 /* numAttrs */, + NULL /* failIdx: unused, we don't need to know which item failed */, + cuda_stream->cuda_stream); +#endif + PARSEC_CUDA_CHECK_ERROR( "cudaMemcpyBatchAsync", cudaStatus, {return PARSEC_ERROR;} ); + return PARSEC_SUCCESS; +} +#endif /* CUDART_VERSION >= 12080 */ + static int parsec_cuda_event_record(struct parsec_device_gpu_module_s *gpu, struct parsec_gpu_exec_stream_s *gpu_stream, int32_t event_idx) { parsec_cuda_exec_stream_t *cuda_stream = (parsec_cuda_exec_stream_t*)gpu_stream; @@ -577,6 +629,11 @@ parsec_cuda_module_init( int dev_id, parsec_device_module_t** module ) device->all_devices_attached = parsec_cuda_all_devices_attached; gpu_device->set_device = parsec_cuda_set_device; gpu_device->memcpy_async = parsec_cuda_memcpy_async; +#if CUDART_VERSION >= 12080 + gpu_device->memcpy_multi_async = parsec_cuda_memcpy_multi_async; +#else + gpu_device->memcpy_multi_async = parsec_device_generic_memcpy_multi_async; +#endif gpu_device->event_record = parsec_cuda_event_record; gpu_device->event_query = parsec_cuda_event_query; gpu_device->memory_info = parsec_cuda_memory_info; diff --git a/parsec/mca/device/device_gpu.c b/parsec/mca/device/device_gpu.c index 67e7ca6c3..5064836b4 100644 --- a/parsec/mca/device/device_gpu.c +++ b/parsec/mca/device/device_gpu.c @@ -1611,9 +1611,30 @@ parsec_device_data_reserve_space( parsec_device_gpu_module_t* gpu_device, return PARSEC_HOOK_RETURN_DONE; } +/** + * @brief Generic fallback for parsec_device_memcpy_multi_async_fn_t: issues @p nb_items + * individual gpu->memcpy_async() calls. Does not stop at the first failure -- every + * item is attempted regardless, and the first error encountered (if any) is returned. + * Used by backends that do not (yet) provide a native batched copy primitive. + */ +int +parsec_device_generic_memcpy_multi_async(parsec_device_gpu_module_t *gpu, parsec_gpu_exec_stream_t *gpu_stream, + void **dsts, void **srcs, size_t *sizes, + parsec_device_transfer_direction_t *directions, int nb_items) +{ + int ret, rc = PARSEC_SUCCESS; + + for(int i = 0; i < nb_items; i++) { + ret = gpu->memcpy_async(gpu, gpu_stream, dsts[i], srcs[i], sizes[i], directions[i]); + if( (PARSEC_SUCCESS == rc) && (PARSEC_SUCCESS != ret) ) + rc = ret; + } + return rc; +} + /* Default stage_in function to transfer data to the GPU device. - * Transfer transfer the contiguous bytes from - * task->data[i].data_in to task->data[i].data_out. + * Transfer the contiguous bytes for every flow set in flow_mask from + * task->data[i].data_in to task->data[i].data_out, as a single multi-item transfer. * * @param[in] task parsec_task_t containing task->data[i].data_in, task->data[i].data_out. * @param[in] flow_mask indicating task flows for which to transfer. @@ -1625,14 +1646,18 @@ parsec_default_gpu_stage_in(parsec_gpu_task_t *gtask, uint32_t flow_mask, parsec_gpu_exec_stream_t *gpu_stream) { - int ret; + int ret, nb_items = 0; parsec_data_copy_t * src_copy; parsec_data_copy_t * dst_copy; parsec_device_gpu_module_t *src_dev; - parsec_device_gpu_module_t *dst_dev; + parsec_device_gpu_module_t *dst_dev = NULL; parsec_task_t *task = gtask->ec; size_t count; parsec_device_transfer_direction_t dir; + void *dsts[MAX_PARAM_COUNT]; + void *srcs[MAX_PARAM_COUNT]; + size_t sizes[MAX_PARAM_COUNT]; + parsec_device_transfer_direction_t directions[MAX_PARAM_COUNT]; for(uint32_t i = 0; i < gtask->nb_flows /* not task->task_class->nb_flows */; i++) { if( !(flow_mask & (1U << i)) ) continue; @@ -1650,20 +1675,24 @@ parsec_default_gpu_stage_in(parsec_gpu_task_t *gtask, } count = (src_copy->original->span <= dst_copy->original->span) ? src_copy->original->span : dst_copy->original->span; - ret = dst_dev->memcpy_async(dst_dev, gpu_stream, - dst_copy->device_private, - src_copy->device_private, - count, - dir); - if(PARSEC_SUCCESS != ret) - return PARSEC_HOOK_RETURN_ERROR; - } + assert(nb_items < MAX_PARAM_COUNT); + dsts[nb_items] = dst_copy->device_private; + srcs[nb_items] = src_copy->device_private; + sizes[nb_items] = count; + directions[nb_items] = dir; + nb_items++; + } + if( 0 == nb_items ) + return PARSEC_HOOK_RETURN_DONE; + ret = dst_dev->memcpy_multi_async( dst_dev, gpu_stream, dsts, srcs, sizes, directions, nb_items ); + if(PARSEC_SUCCESS != ret) + return PARSEC_HOOK_RETURN_ERROR; return PARSEC_HOOK_RETURN_DONE; } /* Default stage_out function to transfer data from the GPU device. - * Transfer transfer the contiguous bytes from - * task->data[i].data_in to task->data[i].data_out. + * Transfer the contiguous bytes for every flow set in flow_mask from + * task->data[i].data_in to task->data[i].data_out, as a single multi-item transfer. * * @param[in] task parsec_task_t containing task->data[i].data_in, task->data[i].data_out. * @param[in] flow_mask indicating task flows for which to transfer. @@ -1675,13 +1704,17 @@ parsec_default_gpu_stage_out(parsec_gpu_task_t *gtask, uint32_t flow_mask, parsec_gpu_exec_stream_t *gpu_stream) { - int ret; + int ret, nb_items = 0; parsec_data_copy_t * src_copy; parsec_data_copy_t * dst_copy; - parsec_device_gpu_module_t *dst_dev, *src_dev; + parsec_device_gpu_module_t *dst_dev, *src_dev = NULL; parsec_task_t *task = gtask->ec; size_t count; parsec_device_transfer_direction_t dir; + void *dsts[MAX_PARAM_COUNT]; + void *srcs[MAX_PARAM_COUNT]; + size_t sizes[MAX_PARAM_COUNT]; + parsec_device_transfer_direction_t directions[MAX_PARAM_COUNT]; for(uint32_t i = 0; i < gtask->nb_flows /* not task->task_class->nb_flows */; i++){ if(flow_mask & (1U << i)){ @@ -1710,16 +1743,20 @@ parsec_default_gpu_stage_out(parsec_gpu_task_t *gtask, return PARSEC_HOOK_RETURN_ERROR; } } - ret = src_dev->memcpy_async( src_dev, gpu_stream, - dst_copy->device_private, - src_copy->device_private, - count, - dir ); - if(PARSEC_SUCCESS != ret) { - return PARSEC_HOOK_RETURN_ERROR; - } + assert(nb_items < MAX_PARAM_COUNT); + dsts[nb_items] = dst_copy->device_private; + srcs[nb_items] = src_copy->device_private; + sizes[nb_items] = count; + directions[nb_items] = dir; + nb_items++; } } + if( 0 == nb_items ) + return PARSEC_HOOK_RETURN_DONE; + ret = src_dev->memcpy_multi_async( src_dev, gpu_stream, dsts, srcs, sizes, directions, nb_items ); + if(PARSEC_SUCCESS != ret) { + return PARSEC_HOOK_RETURN_ERROR; + } return PARSEC_HOOK_RETURN_DONE; } @@ -1790,18 +1827,21 @@ parsec_gpu_data_copy_release_reader(parsec_device_gpu_module_t *gpu_device, /** * If the most current version of the data is not yet available on the GPU memory - * schedule a transfer. + * decide whether a transfer is needed, and if so mark this flow's bit in + * *transfer_mask. The caller is responsible for issuing the actual gpu_task->stage_in + * call (once, covering every flow whose bit got set) after all flows have been decided. * Returns hook special return codes or a positive number: * HOOK_DONE: The most recent version of the data is already available on the GPU - * 1: A copy has been scheduled on the corresponding stream - * HOOK_ERROR: A copy cannot be issued due to GPU. + * 1: A transfer is needed (bit set in *transfer_mask) or already scheduled/pending + * HOOK_AGAIN / HOOK_NEXT: Retry this flow later, no ownership/coherency state changed */ static inline int parsec_device_data_stage_in( parsec_device_gpu_module_t* gpu_device, const parsec_flow_t *flow, parsec_data_pair_t* task_data, parsec_gpu_task_t *gpu_task, - parsec_gpu_exec_stream_t *gpu_stream ) + parsec_gpu_exec_stream_t *gpu_stream, + uint32_t *transfer_mask ) { int32_t type = flow->flow_flags; parsec_data_copy_t *candidate = task_data->data_in; /* best candidate for now */ @@ -2122,24 +2162,13 @@ parsec_device_data_stage_in( parsec_device_gpu_module_t* gpu_device, } #endif gpu_task->flow_info[flow->flow_index].source = candidate; /* save the candidate for release on transfer completion */ - /* Push data into the GPU from the source device */ - int rc = gpu_task->stage_in ? gpu_task->stage_in(gpu_task, (1U << flow->flow_index), gpu_stream): PARSEC_SUCCESS; - if(PARSEC_SUCCESS != rc) { - parsec_warning( "GPU[%d:%s]: gpu_task->stage_in to device rc=%d @%s:%d\n" - "\t<<%p on device %d:%s>> -> <<%p on device %d:%s>> [%zu, %s]", - gpu_device->super.device_index, gpu_device->super.name, rc, __func__, __LINE__, - candidate->device_private, candidate_dev->super.device_index, candidate_dev->super.name, - gpu_elem->device_private, gpu_device->super.device_index, gpu_device->super.name, - span, - (candidate_dev->super.type & gpu_device->super.type & PARSEC_DEV_ANY_TYPE)? "D2D": "H2D"); - if( source_acquired ) { - int readers = parsec_gpu_data_copy_release_reader(candidate_dev, candidate, 1); - assert(readers >= 0); - } - parsec_atomic_unlock( &original->lock ); - assert(0); - return PARSEC_HOOK_RETURN_ERROR; - } + /* Mark this flow as needing a transfer from the source device; the caller issues + * the actual gpu_task->stage_in call once, covering every flow marked this way. + * If that combined call ultimately fails, the caller is responsible for releasing + * any GPU reader acquired above (source_acquired) -- see + * parsec_device_kernel_push_release_readers_on_failure(). + */ + *transfer_mask |= (1U << flow->flow_index); assert(candidate_dev->super.device_index < gpu_device->super.data_in_array_size); gpu_device->super.data_in_from_device[candidate_dev->super.device_index] += span; if( PARSEC_GPU_TASK_TYPE_KERNEL == gpu_task->task_type ) @@ -2730,6 +2759,31 @@ parsec_device_progress_stream( parsec_device_gpu_module_t* gpu_device, goto grab_a_task; } +/** + * @brief Release any GPU reader acquired by parsec_device_data_stage_in() for the + * flows marked in @p transfer_mask. Called when the combined gpu_task->stage_in + * call covering those flows ultimately fails: no completion event will ever fire + * for them, so the normal (transfer-completion) release path will never run. + * A reader was acquired for flow i iff its resolved source is GPU-resident and + * the flow has READ access -- this is externally derivable from + * gpu_task->flow_info[i].source/.flow without needing extra per-flow state. + */ +static inline void +parsec_device_kernel_push_release_readers_on_failure(parsec_gpu_task_t *gpu_task, uint32_t transfer_mask) +{ + for(uint32_t i = 0; i < gpu_task->nb_flows; i++) { + if( !(transfer_mask & (1U << i)) ) continue; + parsec_data_copy_t *src = gpu_task->flow_info[i].source; + const parsec_flow_t *flow = gpu_task->flow_info[i].flow; + if( (NULL == src) || (NULL == flow) ) continue; + if( !(flow->flow_flags & PARSEC_FLOW_ACCESS_READ) ) continue; + parsec_device_module_t *src_dev_mod = parsec_mca_device_get(src->device_index); + if( (NULL == src_dev_mod) || !PARSEC_DEV_IS_GPU(src_dev_mod->type) ) continue; + int readers = parsec_gpu_data_copy_release_reader((parsec_device_gpu_module_t*)src_dev_mod, src, 1); + assert(readers >= 0); + } +} + /** * @brief This function prepare memory on the target device for all the inputs and output * of the task, and then initiate the necessary copies from the best location of the input @@ -2751,6 +2805,7 @@ parsec_device_kernel_push( parsec_device_gpu_module_t *gpu_device, parsec_task_t *this_task = gpu_task->ec; const parsec_flow_t *flow; int ret = 0, input_stream_work = 0; + uint32_t transfer_mask = 0; #if defined(PARSEC_DEBUG_NOISIER) char tmp[MAX_TASK_STRLEN]; #endif @@ -2840,8 +2895,23 @@ parsec_device_kernel_push( parsec_device_gpu_module_t *gpu_device, gpu_device->super.device_index, gpu_device->super.name, flow->name, this_task->data[i].data_out->original->key); ret = parsec_device_data_stage_in( gpu_device, flow, - &(this_task->data[i]), gpu_task, gpu_stream ); + &(this_task->data[i]), gpu_task, gpu_stream, &transfer_mask ); if( ret < 0 ) { + /* Flush any transfers already decided for earlier flows before propagating + * this error/retry: those flows are already marked UNDER_TRANSFER and must + * have a matching physical copy actually issued for them. */ + if( transfer_mask != 0 ) { + int rc = gpu_task->stage_in ? gpu_task->stage_in(gpu_task, transfer_mask, gpu_stream) : PARSEC_SUCCESS; + if( PARSEC_SUCCESS != rc ) { + parsec_warning( "GPU[%d:%s]: gpu_task->stage_in to device rc=%d @%s:%d for task %s transfer_mask=0x%x", + gpu_device->super.device_index, gpu_device->super.name, rc, __func__, __LINE__, + this_task->task_class->name, transfer_mask); + parsec_device_kernel_push_release_readers_on_failure(gpu_task, transfer_mask); + assert(0); + gpu_task->last_status = PARSEC_HOOK_RETURN_ERROR; + return PARSEC_HOOK_RETURN_ERROR; + } + } gpu_task->last_status = ret; return ret; } @@ -2852,6 +2922,18 @@ parsec_device_kernel_push( parsec_device_gpu_module_t *gpu_device, */ input_stream_work += ret; } + if( transfer_mask != 0 ) { + int rc = gpu_task->stage_in ? gpu_task->stage_in(gpu_task, transfer_mask, gpu_stream) : PARSEC_SUCCESS; + if( PARSEC_SUCCESS != rc ) { + parsec_warning( "GPU[%d:%s]: gpu_task->stage_in to device rc=%d @%s:%d for task %s transfer_mask=0x%x", + gpu_device->super.device_index, gpu_device->super.name, rc, __func__, __LINE__, + this_task->task_class->name, transfer_mask); + parsec_device_kernel_push_release_readers_on_failure(gpu_task, transfer_mask); + assert(0); + gpu_task->last_status = PARSEC_HOOK_RETURN_ERROR; + return PARSEC_HOOK_RETURN_ERROR; + } + } PARSEC_DEBUG_VERBOSE(10, parsec_gpu_output_stream, "GPU[%d:%s]: Push task %s DONE", gpu_device->super.device_index, gpu_device->super.name, @@ -2957,6 +3039,8 @@ parsec_device_kernel_pop( parsec_device_gpu_module_t *gpu_device, #endif if (gpu_task->task_type == PARSEC_GPU_TASK_TYPE_D2HTRANSFER) { + uint32_t transfer_mask = 0; + parsec_data_copy_t *cpu_copies[MAX_PARAM_COUNT]; for( int i = 0; i < this_task->locals[0].value; i++ ) { gpu_copy = this_task->data[i].data_out; /* If the gpu copy is not owned by parsec, we don't manage it at all */ @@ -2977,13 +3061,16 @@ parsec_device_kernel_pop( parsec_device_gpu_module_t *gpu_device, goto release_and_return_error; } assert(cpu_copy->data_transfer_status != PARSEC_DATA_STATUS_UNDER_TRANSFER); - rc = gpu_task->stage_out ? gpu_task->stage_out(gpu_task, (1U << i), gpu_stream): PARSEC_SUCCESS; + assert(i < MAX_PARAM_COUNT); + transfer_mask |= (1U << i); + cpu_copies[i] = cpu_copy; + } + if( transfer_mask != 0 ) { + rc = gpu_task->stage_out ? gpu_task->stage_out(gpu_task, transfer_mask, gpu_stream): PARSEC_SUCCESS; if(PARSEC_SUCCESS != rc) { - parsec_warning( "GPU[%d:%s]: gpu_task->stage_out from device rc=%d @%s:%d\n" - "\tdata %s <<%p>> -> <<%p>>\n", + parsec_warning( "GPU[%d:%s]: gpu_task->stage_out from device rc=%d @%s:%d for task %s transfer_mask=0x%x", gpu_device->super.device_index, gpu_device->super.name, rc, __func__, __LINE__, - this_task->task_class->out[i]->name, - gpu_copy->device_private, cpu_copy->device_private); + this_task->task_class->name, transfer_mask); return_code = PARSEC_HOOK_RETURN_DISABLE; goto release_and_return_error; } @@ -2991,8 +3078,11 @@ parsec_device_kernel_pop( parsec_device_gpu_module_t *gpu_device, * owns the copy state transition so custom stage_out callbacks do * not need to know about the GPU copy-transfer bookkeeping. */ - cpu_copy->data_transfer_status = PARSEC_DATA_STATUS_UNDER_TRANSFER; - how_many++; + for( int i = 0; i < this_task->locals[0].value; i++ ) { + if( !(transfer_mask & (1U << i)) ) continue; + cpu_copies[i]->data_transfer_status = PARSEC_DATA_STATUS_UNDER_TRANSFER; + how_many++; + } } return how_many; } @@ -3002,6 +3092,10 @@ parsec_device_kernel_pop( parsec_device_gpu_module_t *gpu_device, gpu_device->super.device_index, gpu_device->super.name, parsec_task_snprintf(tmp, MAX_TASK_STRLEN, this_task) ); + uint32_t transfer_mask = 0; + size_t pending_bytes = 0; + parsec_data_copy_t *pushout_cpu_copies[MAX_PARAM_COUNT]; + for( uint32_t i = 0; i < gpu_task->nb_flows /* not this_task->task_class->nb_flows */; i++ ) { /* We need to manage all data that has been used as input, even if they were read only */ @@ -3133,26 +3227,16 @@ parsec_device_kernel_pop( parsec_device_gpu_module_t *gpu_device, } } #endif - /* Move the data back into main memory */ + /* Defer the actual transfer: accumulate this flow into transfer_mask and + * issue one combined gpu_task->stage_out call after the full loop. The + * cpu_copy state transition is likewise deferred to a second pass over + * the masked flows, taken under each flow's own original->lock, once the + * combined call has actually succeeded (see below). */ assert(cpu_copy->data_transfer_status != PARSEC_DATA_STATUS_UNDER_TRANSFER); - rc = gpu_task->stage_out? gpu_task->stage_out(gpu_task, (1U << flow->flow_index), gpu_stream): PARSEC_SUCCESS; - if( PARSEC_SUCCESS != rc ) { - parsec_warning( "GPU[%d:%s]: gpu_task->stage_out from device rc=%d @%s:%d\n" - "\tdata %s <<%p>> -> <<%p>>\n", - gpu_device->super.device_index, gpu_device->super.name, rc, __func__, __LINE__, - this_task->task_class->out[i]->name, - gpu_copy->device_private, cpu_copy->device_private); - return_code = PARSEC_HOOK_RETURN_DISABLE; - parsec_atomic_unlock(&original->lock); - goto release_and_return_error; - } - /* stage_out only enqueues the device-to-host transfer. The runtime - * owns the copy state transition so custom stage_out callbacks do - * not need to know about the GPU copy-transfer bookkeeping. - */ - cpu_copy->data_transfer_status = PARSEC_DATA_STATUS_UNDER_TRANSFER; - gpu_device->super.data_out_to_host += span; /* TODO: not hardcoded, use datatype size */ - how_many++; + assert(flow->flow_index < MAX_PARAM_COUNT); + transfer_mask |= (1U << flow->flow_index); + pushout_cpu_copies[flow->flow_index] = cpu_copy; + pending_bytes += span; /* TODO: not hardcoded, use datatype size */ } else { assert( 0 == gpu_copy->readers ); } @@ -3160,6 +3244,32 @@ parsec_device_kernel_pop( parsec_device_gpu_module_t *gpu_device, parsec_atomic_unlock(&original->lock); } + if( transfer_mask != 0 ) { + rc = gpu_task->stage_out? gpu_task->stage_out(gpu_task, transfer_mask, gpu_stream): PARSEC_SUCCESS; + if(PARSEC_SUCCESS != rc) { + parsec_warning( "GPU[%d:%s]: gpu_task->stage_out from device rc=%d @%s:%d for task %s transfer_mask=0x%x", + gpu_device->super.device_index, gpu_device->super.name, rc, __func__, __LINE__, + this_task->task_class->name, transfer_mask); + return_code = PARSEC_HOOK_RETURN_DISABLE; + goto release_and_return_error; + } + /* stage_out only enqueues the device-to-host transfer. The runtime owns the + * copy state transition so custom stage_out callbacks do not need to know + * about the GPU copy-transfer bookkeeping. Re-acquire each flow's original + * lock briefly, matching the protection this state transition had when it + * used to happen inline under that same lock. */ + for( uint32_t i = 0; i < gpu_task->nb_flows; i++ ) { + if( !(transfer_mask & (1U << i)) ) continue; + parsec_data_copy_t *cpu_copy = pushout_cpu_copies[i]; + parsec_data_t *cpu_copy_original = cpu_copy->original; + parsec_atomic_lock(&cpu_copy_original->lock); + cpu_copy->data_transfer_status = PARSEC_DATA_STATUS_UNDER_TRANSFER; + parsec_atomic_unlock(&cpu_copy_original->lock); + how_many++; + } + gpu_device->super.data_out_to_host += pending_bytes; + } + release_and_return_error: if( update_data_epoch ) { gpu_device->data_avail_epoch++; diff --git a/parsec/mca/device/device_gpu.h b/parsec/mca/device/device_gpu.h index b36a40718..636486066 100644 --- a/parsec/mca/device/device_gpu.h +++ b/parsec/mca/device/device_gpu.h @@ -183,6 +183,32 @@ typedef int (*parsec_device_set_device_fn_t)(struct parsec_device_gpu_module_s * typedef int (*parsec_device_memcpy_async_fn_t)(struct parsec_device_gpu_module_s *gpu, struct parsec_gpu_exec_stream_s *gpu_stream, void *dest, void *source, size_t bytes, parsec_device_transfer_direction_t direction); +/** + * @brief Schedules the asynchronous copy of @p nb_items entries onto the GPU stream + * of @p gpu_stream, as a single (ideally batched) operation. @p dsts, @p srcs, + * @p sizes, and @p directions are parallel arrays of length @p nb_items: + * entry i copies @p sizes[i] bytes from @p srcs[i] to @p dsts[i], transferring + * across the memory spaces described by @p directions[i]. + * + * The array shape deliberately mirrors native batched copy primitives (e.g. + * cudaMemcpyBatchAsync's dsts/srcs/sizes/count parameters) so that a backend + * with a native batched primitive can forward these arrays directly, without + * repacking them into some other layout first. @p directions is not needed by + * such primitives (the copy kind is inferred from the pointers), but is kept + * so that parsec_device_generic_memcpy_multi_async() -- the fallback used by + * backends without a native batched primitive -- can still issue the correct + * per-item memcpy_async() call. + * + * @details typically maps to a batched copy primitive (e.g. cudaMemcpyBatchAsync) + * if the backend has one, or falls back to issuing @p nb_items individual + * memcpy_async calls (see parsec_device_generic_memcpy_multi_async()). + * + * @return PARSEC_SUCCESS or a PARSEC error + */ +typedef int (*parsec_device_memcpy_multi_async_fn_t)(struct parsec_device_gpu_module_s *gpu, struct parsec_gpu_exec_stream_s *gpu_stream, + void **dsts, void **srcs, size_t *sizes, + parsec_device_transfer_direction_t *directions, int nb_items); + /** * @brief Record an event on the GPU @p gpu_stream of GPU @p gpu, with index @p idx. * @@ -249,6 +275,7 @@ struct parsec_device_gpu_module_s { /* This set of base functions is used by the GPU devices to implement their Device Management Functions */ parsec_device_set_device_fn_t set_device; parsec_device_memcpy_async_fn_t memcpy_async; + parsec_device_memcpy_multi_async_fn_t memcpy_multi_async; parsec_device_event_query_fn_t event_query; parsec_device_event_record_fn_t event_record; parsec_device_memory_info_fn_t memory_info; @@ -419,6 +446,19 @@ parsec_device_kernel_scheduler( parsec_device_module_t *module, parsec_execution_stream_t *es, void *gpu_task ); +/** + * @brief Generic fallback for parsec_device_memcpy_multi_async_fn_t: issues @p nb_items + * individual gpu->memcpy_async() calls. Does not stop at the first failure -- every + * item is attempted regardless, and the first error encountered (if any) is returned. + * Used by backends that do not (yet) provide a native batched copy primitive. + * + * @return PARSEC_SUCCESS or the first PARSEC error encountered + */ +int +parsec_device_generic_memcpy_multi_async(parsec_device_gpu_module_t *gpu, parsec_gpu_exec_stream_t *gpu_stream, + void **dsts, void **srcs, size_t *sizes, + parsec_device_transfer_direction_t *directions, int nb_items); + /* Default stage_in function to transfer data to the GPU device. * Transfer transfer the contiguous bytes from * task->data[i].data_in to task->data[i].data_out. diff --git a/parsec/mca/device/level_zero/device_level_zero_module.c b/parsec/mca/device/level_zero/device_level_zero_module.c index 4b2ef0799..37f56d4c4 100644 --- a/parsec/mca/device/level_zero/device_level_zero_module.c +++ b/parsec/mca/device/level_zero/device_level_zero_module.c @@ -422,6 +422,7 @@ int parsec_level_zero_module_init( int dev_id, parsec_device_level_zero_driver_t device->memory_unregister = NULL; // TODO there seem to be no memory pinning in level_zero? gpu_device->set_device = parsec_level_zero_set_device; gpu_device->memcpy_async = parsec_level_zero_memcpy_async; + gpu_device->memcpy_multi_async = parsec_device_generic_memcpy_multi_async; gpu_device->event_record = parsec_level_zero_event_record; gpu_device->event_query = parsec_level_zero_event_query; gpu_device->memory_info = parsec_level_zero_memory_info;