Skip to content

Commit 0febdba

Browse files
abouteillerabouteiller
committed
Merge remote-tracking branch 'origin/master' into topic/derive_task_from_parsec_object
2 parents de821b8 + c67a2e8 commit 0febdba

5 files changed

Lines changed: 81 additions & 51 deletions

File tree

‎parsec/mca/device/cuda/device_cuda_module.c‎

Lines changed: 24 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
* Copyright (c) 2010-2024 The University of Tennessee and The University
2+
* Copyright (c) 2010-2025 The University of Tennessee and The University
33
* of Tennessee Research Foundation. All rights
44
* reserved.
55
* Copyright (c) 2024 NVIDIA Corporation. All rights reserved.
@@ -307,18 +307,18 @@ static void* parsec_cuda_find_incarnation(parsec_device_gpu_module_t* gpu_device
307307

308308
static int parsec_cuda_set_device(parsec_device_gpu_module_t *gpu)
309309
{
310-
cudaError_t cudaStatus;
310+
cudaError_t status;
311311
parsec_device_cuda_module_t *cuda_device = (parsec_device_cuda_module_t *)gpu;
312312

313-
cudaStatus = cudaSetDevice(cuda_device->cuda_index);
314-
PARSEC_CUDA_CHECK_ERROR( "cudaSetDevice", cudaStatus, {return PARSEC_ERROR;} );
313+
status = cudaSetDevice(cuda_device->cuda_index);
314+
PARSEC_CUDA_CHECK_ERROR( "cudaSetDevice", status, {return PARSEC_ERROR;} );
315315
return PARSEC_SUCCESS;
316316
}
317317

318318
static int parsec_cuda_memcpy_async(struct parsec_device_gpu_module_s *gpu, struct parsec_gpu_exec_stream_s *gpu_stream,
319319
void *dest, void *source, size_t bytes, parsec_device_transfer_direction_t direction)
320320
{
321-
cudaError_t cudaStatus;
321+
cudaError_t status;
322322
enum cudaMemcpyKind kind;
323323
parsec_cuda_exec_stream_t *cuda_stream = (parsec_cuda_exec_stream_t *)gpu_stream;
324324

@@ -338,45 +338,45 @@ static int parsec_cuda_memcpy_async(struct parsec_device_gpu_module_s *gpu, stru
338338
PARSEC_CUDA_CHECK_ERROR( "Translate parsec_device_transfer_direction_t to cudaMemcpyKind", cudaErrorInvalidValue, {return PARSEC_ERROR;} );
339339
}
340340

341-
cudaStatus = cudaMemcpyAsync( dest, source, bytes, kind, cuda_stream->cuda_stream );
342-
PARSEC_CUDA_CHECK_ERROR( "cudaMemcpyAsync", cudaStatus, {return PARSEC_ERROR;} );
341+
status = cudaMemcpyAsync( dest, source, bytes, kind, cuda_stream->cuda_stream );
342+
PARSEC_CUDA_CHECK_ERROR( "cudaMemcpyAsync", status, {return PARSEC_ERROR;} );
343343
return PARSEC_SUCCESS;
344344
}
345345

346346
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)
347347
{
348348
parsec_cuda_exec_stream_t *cuda_stream = (parsec_cuda_exec_stream_t*)gpu_stream;
349-
cudaError_t cudaStatus;
349+
cudaError_t status;
350350
(void)gpu;
351351

352-
cudaStatus = cudaEventRecord(cuda_stream->events[event_idx], cuda_stream->cuda_stream);
353-
PARSEC_CUDA_CHECK_ERROR( "cudaEventRecord", cudaStatus, {return PARSEC_ERROR;} );
352+
status = cudaEventRecord(cuda_stream->events[event_idx], cuda_stream->cuda_stream);
353+
PARSEC_CUDA_CHECK_ERROR( "cudaEventRecord", status, {return PARSEC_ERROR;} );
354354
return PARSEC_SUCCESS;
355355
}
356356

357357
static int parsec_cuda_event_query(struct parsec_device_gpu_module_s *gpu, struct parsec_gpu_exec_stream_s *gpu_stream, int32_t event_idx)
358358
{
359359
parsec_cuda_exec_stream_t *cuda_stream = (parsec_cuda_exec_stream_t*)gpu_stream;
360-
cudaError_t cudaStatus;
360+
cudaError_t status;
361361
(void)gpu;
362362

363-
cudaStatus = cudaEventQuery(cuda_stream->events[event_idx]);
364-
if(cudaSuccess == cudaStatus) {
363+
status = cudaEventQuery(cuda_stream->events[event_idx]);
364+
if(cudaSuccess == status) {
365365
return 1;
366366
}
367-
if(cudaErrorNotReady == cudaStatus) {
367+
if(cudaErrorNotReady == status) {
368368
return 0;
369369
}
370-
PARSEC_CUDA_CHECK_ERROR( "cudaEventQuery", cudaStatus, {return PARSEC_ERROR;} );
370+
PARSEC_CUDA_CHECK_ERROR( "cudaEventQuery", status, {return PARSEC_ERROR;} );
371371
return PARSEC_ERROR; /* should be unreachable */
372372
}
373373

374374
static int parsec_cuda_memory_info(struct parsec_device_gpu_module_s *gpu, size_t *free_mem, size_t *total_mem)
375375
{
376-
cudaError_t cudaStatus;
376+
cudaError_t status;
377377
(void)gpu;
378-
cudaStatus = cudaMemGetInfo(free_mem, total_mem);
379-
PARSEC_CUDA_CHECK_ERROR( "cudaMemGetInfo", cudaStatus, {return PARSEC_ERROR;});
378+
status = cudaMemGetInfo(free_mem, total_mem);
379+
PARSEC_CUDA_CHECK_ERROR( "cudaMemGetInfo", status, {return PARSEC_ERROR;});
380380
/* sanity check: no more free memory than total memory
381381
* some hip devices on Frontier may report more free memory than is available */
382382
if (*free_mem > *total_mem) {
@@ -387,19 +387,19 @@ static int parsec_cuda_memory_info(struct parsec_device_gpu_module_s *gpu, size_
387387

388388
static int parsec_cuda_memory_allocate(struct parsec_device_gpu_module_s *gpu, size_t bytes, void **addr)
389389
{
390-
cudaError_t cudaStatus;
390+
cudaError_t status;
391391
(void)gpu;
392-
cudaStatus = cudaMalloc(addr, bytes);
393-
PARSEC_CUDA_CHECK_ERROR( "cudaMalloc", cudaStatus, {return PARSEC_ERROR;});
392+
status = cudaMalloc(addr, bytes);
393+
PARSEC_CUDA_CHECK_ERROR( "cudaMalloc", status, {return PARSEC_ERROR;});
394394
return PARSEC_SUCCESS;
395395
}
396396

397397
static int parsec_cuda_memory_free(struct parsec_device_gpu_module_s *gpu, void *addr)
398398
{
399-
cudaError_t cudaStatus;
399+
cudaError_t status;
400400
(void)gpu;
401-
cudaStatus = cudaFree(addr);
402-
PARSEC_CUDA_CHECK_ERROR( "cudaFree", cudaStatus, {return PARSEC_ERROR;});
401+
status = cudaFree(addr);
402+
PARSEC_CUDA_CHECK_ERROR( "cudaFree", status, {return PARSEC_ERROR;});
403403
return PARSEC_SUCCESS;
404404
}
405405

‎parsec/mca/device/device_gpu.c‎

Lines changed: 24 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -927,11 +927,31 @@ parsec_device_data_reserve_space( parsec_device_gpu_module_t* gpu_device,
927927
flow->name, i, gpu_elem,
928928
gpu_elem->data_transfer_status == PARSEC_DATA_STATUS_UNDER_TRANSFER ? " [in transfer]" : "");
929929
if ( gpu_elem->data_transfer_status == PARSEC_DATA_STATUS_UNDER_TRANSFER ) {
930-
/* The data is indeed under transfer, but as we always force an event at the end of this
931-
* step, we do not need to have a special case for this, because the forced event will
932-
* ensure the data will be available on the GPU by the time this task will move to the
933-
* next step.
930+
/* The data is under transfer, which is fine for RO data since we always force an event
931+
* at the end of this step so we do not need to have a special case for this. The forced
932+
* event will ensure the data will be available on the GPU by the time this task will move
933+
* to the next step. For WRITE flows, we have to abort here and come back later because
934+
* transfer_ownership will bark at WRITE flows that are under transfer.
934935
*/
936+
if (0 != (PARSEC_FLOW_ACCESS_WRITE & flow->flow_flags)) {
937+
this_task->data[i].data_out = NULL;
938+
PARSEC_DEBUG_VERBOSE(20, parsec_gpu_output_stream,
939+
"GPU[%d:%s]:%s:\tWrite flow under transfer for copy %p",
940+
gpu_device->super.device_index, gpu_device->super.name, task_name,
941+
gpu_elem);
942+
for (j = 0; j <= i; ++j) {
943+
if (temp_loc[i] == NULL) continue;
944+
if( 0 == (temp_loc[j]->flags & PARSEC_DATA_FLAG_PARSEC_OWNED) ) continue;
945+
PARSEC_DEBUG_VERBOSE(20, parsec_gpu_output_stream,
946+
"GPU[%d:%s]:%s:\tAdd copy %p [ref_count %d] back to the LRU list",
947+
gpu_device->super.device_index, gpu_device->super.name, task_name,
948+
temp_loc[j], temp_loc[j]->super.super.obj_reference_count);
949+
parsec_list_push_front(&gpu_device->gpu_mem_lru, (parsec_list_item_t*)temp_loc[j]);
950+
}
951+
parsec_atomic_unlock(&master->lock);
952+
return PARSEC_HOOK_RETURN_AGAIN;
953+
}
954+
935955
}
936956
parsec_atomic_unlock(&master->lock);
937957
continue;

‎parsec/mca/device/hip/ValidateModule.CMake‎

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,10 @@ if( PARSEC_HAVE_HIP )
66
file(GLOB cuda_sources ${MCA_BASE_DIR}/${COMPONENT}/cuda/[^\\.]*.[ch])
77
find_package(Perl REQUIRED)
88
find_program(HIPIFY_PERL_COMMAND NAMES hipify-perl HINTS ${HIP_BIN_INSTALL_DIR} REQUIRED)
9+
option(PARSEC_HIPIFY_MCA_PRINT_STATS "Print the statistics during the CUDA device MCA conversion to HIP" OFF)
10+
if(PARSEC_HIPIFY_MCA_PRINT_STATS)
11+
set(MCA_HIPIFY_FLAGS "--print-stats")
12+
endif(PARSEC_HIPIFY_MCA_PRINT_STATS)
913
set(module_dir "mca/${COMPONENT}/${MODULE}")
1014
foreach(cuda_file ${cuda_sources})
1115
file(RELATIVE_PATH cuda_filename ${MCA_BASE_DIR}/${COMPONENT}/cuda ${cuda_file})
@@ -14,7 +18,7 @@ if( PARSEC_HAVE_HIP )
1418
add_custom_command(OUTPUT ${${MODULE}_file}
1519
DEPENDS ${cuda_file} # do not use MAIN_DEPENDENCY, that overides the default .c.o rule
1620
COMMAND ${CMAKE_COMMAND} -E copy "${cuda_file}" "${${MODULE}_file}.prehip"
17-
COMMAND ${PERL_EXECUTABLE} ${HIPIFY_PERL_COMMAND} --inplace --print-stats "${${MODULE}_file}"
21+
COMMAND ${PERL_EXECUTABLE} ${HIPIFY_PERL_COMMAND} --inplace ${MCA_HIPIFY_FLAGS} "${${MODULE}_file}"
1822
COMMAND ${PERL_EXECUTABLE} -i -pe "s{(cuda)}{ substr uc hip | (uc \$1 ^ \$1), 0, 3 }egi" "${${MODULE}_file}" VERBATIM) # Convert all remaining cuda/CUDA
1923
if(${${MODULE}_file} MATCHES [^\\.]*.c) # do not add .h to sources
2024
list(APPEND MCA_${COMPONENT}_${MODULE}_SOURCES ${${MODULE}_file})

‎parsec/parsec.c‎

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -149,8 +149,8 @@ static char *parsec_dot_file = NULL;
149149
static char *parsec_app_name = NULL;
150150

151151
static int parsec_runtime_max_number_of_cores = -1;
152-
static int parsec_runtime_bind_main_thread = 1;
153-
static int parsec_runtime_bind_threads = 1;
152+
static int parsec_runtime_bind_main_thread = 0;
153+
static int parsec_runtime_bind_threads = 0;
154154

155155
int parsec_runtime_keep_highest_priority_task = 1;
156156

@@ -521,10 +521,11 @@ parsec_context_t* parsec_init( int nb_cores, int* pargc, char** pargv[] )
521521
false, false, vpmap_parameter, &vpmap_parameter);
522522

523523
/* thread binding */
524-
parsec_mca_param_reg_int_name("bind", "main_thread", "Force the binding of the thread calling parsec_init",
525-
false, false, parsec_runtime_bind_main_thread, &parsec_runtime_bind_main_thread);
526524
parsec_mca_param_reg_int_name("bind", "threads", "Bind main and worker threads", false, false,
527525
parsec_runtime_bind_threads, &parsec_runtime_bind_threads);
526+
parsec_runtime_bind_main_thread = parsec_runtime_bind_threads;
527+
parsec_mca_param_reg_int_name("bind", "main_thread", "Force the binding of the thread calling parsec_init",
528+
false, false, parsec_runtime_bind_main_thread, &parsec_runtime_bind_main_thread);
528529
parsec_mca_param_reg_int_name("bind", "comm", "Bind the communication thread to physical core <int>."
529530
"-1: do not bind."
530531
"Warning: binding relies on HWLOC, be careful when using with cgroups",

‎parsec/scheduling.c‎

Lines changed: 23 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -2,6 +2,7 @@
22
* Copyright (c) 2009-2024 The University of Tennessee and The University
33
* of Tennessee Research Foundation. All rights
44
* reserved.
5+
* Copyright (c) 2025 NVIDIA Corporation. All rights reserved.
56
*/
67

78
#include "parsec/runtime.h"
@@ -323,26 +324,22 @@ __parsec_schedule(parsec_execution_stream_t* es,
323324
* stream should never be used concurrently in two call to this function (or
324325
* a thread should never `borrow` an execution stream for this call).
325326
*/
326-
int __parsec_schedule_vp(parsec_execution_stream_t* es,
327+
int __parsec_schedule_vp(parsec_execution_stream_t* submission_es,
327328
parsec_task_t** task_rings,
328329
int32_t distance)
329330
{
330-
parsec_execution_stream_t* target_es;
331-
const parsec_vp_t** vps = (const parsec_vp_t**)es->virtual_process->parsec_context->virtual_processes;
331+
parsec_execution_stream_t *target_es,
332+
*es = (NULL == submission_es ? parsec_my_execution_stream() : submission_es);
332333
int ret = 0;
333334

334-
#if defined(PARSEC_DEBUG_PARANOID)
335-
/* As the setting of the next_task is not protected no thread should call
336-
* this function with a stream other than its own. */
337-
assert( (NULL == es) || (parsec_my_execution_stream() == es) );
338-
#endif /* defined(PARSEC_DEBUG_PARANOID) */
339-
340-
if( NULL == es || !parsec_runtime_keep_highest_priority_task || NULL == es->scheduler_object) {
341-
for(int vp = 0; vp < es->virtual_process->parsec_context->nb_vp; vp++ ) {
335+
if( NULL == submission_es || !parsec_runtime_keep_highest_priority_task ||
336+
distance || NULL == submission_es->scheduler_object /* not the communication thread */ ) {
337+
parsec_context_t* context = es->virtual_process->parsec_context;
338+
for(int vp = 0; vp < context->nb_vp; vp++ ) {
342339
parsec_task_t* ring = task_rings[vp];
343340
if( NULL == ring ) continue;
344341

345-
target_es = vps[vp]->execution_streams[0];
342+
target_es = context->virtual_processes[vp]->execution_streams[0];
346343

347344
ret = __parsec_schedule(target_es, ring, distance);
348345
if( 0 != ret )
@@ -352,22 +349,30 @@ int __parsec_schedule_vp(parsec_execution_stream_t* es,
352349
}
353350
return ret;
354351
}
355-
for(int vp = 0; vp < es->virtual_process->parsec_context->nb_vp; vp++ ) {
352+
#if defined(PARSEC_DEBUG_PARANOID)
353+
/* As the setting of the next_task is not protected no thread should call
354+
* this function with a stream other than its own (with a distance 0). */
355+
assert( (parsec_my_execution_stream() == submission_es) );
356+
#endif /* defined(PARSEC_DEBUG_PARANOID) */
357+
358+
parsec_context_t* context = submission_es->virtual_process->parsec_context;
359+
for(int vp = 0; vp < context->nb_vp; vp++ ) {
356360
parsec_task_t* ring = task_rings[vp];
357361
if( NULL == ring ) continue;
358362

359-
target_es = vps[vp]->execution_streams[0];
363+
target_es = context->virtual_processes[vp]->execution_streams[0];
360364

361-
if( vp == es->virtual_process->vp_id ) {
362-
if( NULL == es->next_task ) {
363-
es->next_task = ring;
365+
if( vp == submission_es->virtual_process->vp_id ) {
366+
if( NULL == submission_es->next_task ) {
367+
submission_es->next_task = ring;
364368
ring = (parsec_task_t*)parsec_list_item_ring_chop(&ring->super);
365369
if( NULL == ring ) {
366370
task_rings[vp] = NULL; /* remove the tasks already scheduled */
367371
continue;
368372
}
369373
}
370-
target_es = es;
374+
/* Beware we are changing the submission execution stream for the local vp */
375+
target_es = submission_es;
371376
}
372377
ret = __parsec_schedule(target_es, ring, distance);
373378
if( 0 != ret )

0 commit comments

Comments
 (0)