From f70b109cc37cf2d38fd1ea4e343331257642eced Mon Sep 17 00:00:00 2001 From: Tao Sang Date: Mon, 14 Jun 2021 19:34:26 -0400 Subject: [PATCH] SWDEV-291009 - Fix LaunchCoopMultiKernel on NV Fix compiling error of hipLaunchCoopMultiKernel on NV A100. Fix test failure on all NV devices. Change-Id: Iba20caa0e9021480378625506197384c275a3289 --- .../cooperativeGrps/hipLaunchCoopMultiKernel.cpp | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/tests/src/runtimeApi/cooperativeGrps/hipLaunchCoopMultiKernel.cpp b/tests/src/runtimeApi/cooperativeGrps/hipLaunchCoopMultiKernel.cpp index 171838eb11..239d552708 100644 --- a/tests/src/runtimeApi/cooperativeGrps/hipLaunchCoopMultiKernel.cpp +++ b/tests/src/runtimeApi/cooperativeGrps/hipLaunchCoopMultiKernel.cpp @@ -20,7 +20,7 @@ THE SOFTWARE. // Simple test for hipLaunchCooperativeKernelMultiDevice API. /* HIT_START - * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 -rdc=true -gencode arch=compute_70,code=sm_70 + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 -rdc=true -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80 * TEST: %t * HIT_END */ @@ -123,7 +123,7 @@ int main() { HIPCHECK(hipMalloc((void**)&dA[i], SIZE)); HIPCHECK(hipMalloc((void**)&dB[i], 64 * deviceProp[i].multiProcessorCount * sizeof(long))); if (i == 0) { - HIPCHECK(hipHostMalloc((void**)&dC, (nGpu + 1) * sizeof(long), hipHostMallocCoherent)); + HIPCHECK(hipHostMalloc((void**)&dC, (nGpu + 1) * sizeof(long))); } HIPCHECK(hipMemcpy(dA[i], &init[i * copySizeInDwords] , SIZE, hipMemcpyHostToDevice)); HIPCHECK(hipStreamCreate(&stream[i])); @@ -175,6 +175,9 @@ int main() { system_clock::time_point start = system_clock::now(); hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0); + for (int i = 0; i < nGpu; i++) { + hipStreamSynchronize(stream[i]); + } system_clock::time_point end = system_clock::now(); std::chrono::duration elapsed_seconds = end - start; end_time = std::chrono::system_clock::to_time_t(end);