From 1ae1c3ea44ffae3d9f00d421607e9f5948f52fca Mon Sep 17 00:00:00 2001 From: gobhardw Date: Mon, 26 Jun 2023 23:24:20 +0530 Subject: [PATCH] Fixing copy_on_engine app Change-Id: Icd53a9eead2d8ca422838cd2843860c9ead57fff [ROCm/rocprofiler commit: f1b097e37e5af0eebb6740697e872803157fccd5] --- .../featuretests/tracer/apps/copy_on_engine.cpp | 14 +++++++++----- 1 file changed, 9 insertions(+), 5 deletions(-) diff --git a/projects/rocprofiler/tests-v2/featuretests/tracer/apps/copy_on_engine.cpp b/projects/rocprofiler/tests-v2/featuretests/tracer/apps/copy_on_engine.cpp index ec5d89b578..8738067c9d 100644 --- a/projects/rocprofiler/tests-v2/featuretests/tracer/apps/copy_on_engine.cpp +++ b/projects/rocprofiler/tests-v2/featuretests/tracer/apps/copy_on_engine.cpp @@ -208,16 +208,20 @@ static hsa_status_t AsyncCpyTest(async_mem_cpy_agent* dst, async_mem_cpy_agent* err = hsa_signal_create(1, 0, NULL, ©_signal); RET_IF_HSA_ERR(err); // Do the copy... - err = hsa_amd_memory_async_copy(dst->ptr, dst->dev, src->ptr, src->dev, sz, 0, NULL, copy_signal); - RET_IF_HSA_ERR(err); + // err = hsa_amd_memory_async_copy(dst->ptr, dst->dev, src->ptr, src->dev, sz, 0, NULL, + // copy_signal); RET_IF_HSA_ERR(err); // call following APIs to make sure we intercept hsa_amd_memory_async_copy_on_engine uint32_t engine_ids_mask = 0; err = hsa_amd_memory_copy_engine_status(args->cpu.dev, args->gpu1.dev, &engine_ids_mask); - hsa_amd_sdma_engine_id_t engine_id = HSA_AMD_SDMA_ENGINE_0; - err = hsa_amd_memory_async_copy_on_engine(dst->ptr, dst->dev, src->ptr, src->dev, sz, 0, NULL, - copy_signal, engine_id, false); + RET_IF_HSA_ERR(err); + uint32_t engine_id = HSA_AMD_SDMA_ENGINE_0 & engine_ids_mask; + err = hsa_amd_memory_async_copy_on_engine( + dst->ptr, dst->dev, src->ptr, src->dev, sz, 0, NULL, copy_signal, + static_cast(engine_id), false); + RET_IF_HSA_ERR(err); + // Here we do a blocking wait. Alternatively, we could also use a // non-blocking wait in a loop, and do other work while waiting. if (hsa_signal_wait_relaxed(copy_signal, HSA_SIGNAL_CONDITION_LT, 1, -1,