From 9e2e331f786aa9433eae9af4f9656087869bde2a Mon Sep 17 00:00:00 2001 From: Sourabh Betigeri Date: Thu, 26 Sep 2024 17:31:47 -0700 Subject: [PATCH] SWDEV-450052 - Fixes for negative cases of hipExtLaunchMultiKernelMultiDevice No known documentation suggesting launchParamsList.func/gridDim/blockDim/sharedMem should not match across all devices. Since this is a Ext AMD only API, these can be a AMD-only feature of this API. Change-Id: Icbea4eda11ec936d2a76fdcb44bfd9441877adf6 --- .../hipExtLaunchMultiKernelMultiDevice.cc | 28 +------------------ 1 file changed, 1 insertion(+), 27 deletions(-) diff --git a/catch/unit/executionControl/hipExtLaunchMultiKernelMultiDevice.cc b/catch/unit/executionControl/hipExtLaunchMultiKernelMultiDevice.cc index 97b1420b9..51bc64bd2 100644 --- a/catch/unit/executionControl/hipExtLaunchMultiKernelMultiDevice.cc +++ b/catch/unit/executionControl/hipExtLaunchMultiKernelMultiDevice.cc @@ -90,32 +90,6 @@ TEST_CASE("Unit_hipExtLaunchMultiKernelMultiDevice_Negative_Parameters") { hipErrorInvalidValue); } - if (device_count > 1) { - SECTION("launchParamsList.func doesn't match across all devices") { - params_list[1].func = reinterpret_cast(kernel2); - HIP_CHECK_ERROR(hipExtLaunchMultiKernelMultiDevice(params_list.data(), device_count, 0u), - hipErrorInvalidValue); - } - - SECTION("launchParamsList.gridDim doesn't match across all kernels") { - params_list[1].gridDim = dim3{2, 2, 2}; - HIP_CHECK_ERROR(hipExtLaunchMultiKernelMultiDevice(params_list.data(), device_count, 0u), - hipErrorInvalidValue); - } - - SECTION("launchParamsList.blockDim doesn't match across all kernels") { - params_list[1].blockDim = dim3{2, 2, 2}; - HIP_CHECK_ERROR(hipExtLaunchMultiKernelMultiDevice(params_list.data(), device_count, 0u), - hipErrorInvalidValue); - } - - SECTION("launchParamsList.sharedMem doesn't match across all kernels") { - params_list[1].sharedMem = 1024; - HIP_CHECK_ERROR(hipExtLaunchMultiKernelMultiDevice(params_list.data(), device_count, 0u), - hipErrorInvalidValue); - } - } - for (const auto params : params_list) { HIP_CHECK(hipStreamDestroy(params.stream)); } @@ -136,7 +110,7 @@ TEST_CASE("Unit_hipExtLaunchMultiKernelMultiDevice_Negative_MultiKernelSameDevic } HIP_CHECK_ERROR(hipExtLaunchMultiKernelMultiDevice(params_list.data(), 2, 0u), - hipErrorInvalidValue); + hipErrorInvalidDevice); for (const auto params : params_list) { HIP_CHECK(hipStreamDestroy(params.stream));