EXSWHTEC-114 - Implement tests for hipExtLaunchKernel APIs #60

Change-Id: I9152c31c2c07ddbfb48865cd68a42557e763be28 [ROCm/hip-tests commit: 3f65d7a410]
2023-12-28 19:32:13 +01:00
@@ -53,6 +53,10 @@
        "Unit_hipFuncSetSharedMemConfig_Negative_Not_Supported",
        "Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize_Not_Supported",
        "Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout_Not_Supported",
+        "NOTE: The following test is disabled due to defect - EXSWHTEC-243",
+        "Unit_hipExtLaunchKernel_Negative_Parameters",
+        "NOTE: The following test is disabled due to defect - EXSWHTEC-244",
+        "Unit_hipExtLaunchMultiKernelMultiDevice_Negative_Parameters",
        "Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters",
        "Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Negative_Parameters",
        "Unit_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_Negative_Parameters",
@@ -218,6 +218,10 @@
        "Unit_hipVectorTypes_test_on_device",
        "=== Patch which removes the typetraits implementation from std namespace in hiprtc is reverted ===",
        "Unit_hiprtc_stdheaders",
+        "NOTE: The following test is disabled due to defect - EXSWHTEC-243",
+        "Unit_hipExtLaunchKernel_Negative_Parameters",
+        "NOTE: The following test is disabled due to defect - EXSWHTEC-244",
+        "Unit_hipExtLaunchMultiKernelMultiDevice_Negative_Parameters",
        "Unit_hipMemAddressReserve_AlignmentTest",
        "Unit_hipGraphAddMemcpyNode_Negative_Parameters",
        "Unit_hipMemCreate_ChkWithKerLaunch",
@@ -5,6 +5,13 @@ set(TEST_SRC
    hipFuncSetAttribute.cc
 )

+if(HIP_PLATFORM MATCHES "amd")
+    set(TEST_SRC ${TEST_SRC}
+        hipExtLaunchKernel.cc
+        hipExtLaunchMultiKernelMultiDevice.cc
+    )
+endif()
+
 hip_add_exe_to_target(NAME ExecutionControlTest
                      TEST_SRC ${TEST_SRC}
                      TEST_TARGET_NAME build_tests
@@ -0,0 +1,176 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "execution_control_common.hh"
+
+#include <hip_test_common.hh>
+#include <hip/hip_runtime_api.h>
+#include <resource_guards.hh>
+#include <utils.hh>
+
+TEST_CASE("Unit_hipExtLaunchKernel_Positive_Basic") {
+  SECTION("Kernel with no arguments") {
+    HIP_CHECK(hipExtLaunchKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1}, dim3{1, 1, 1},
+                                 nullptr, 0, nullptr, nullptr, nullptr, 0u));
+    HIP_CHECK(hipDeviceSynchronize());
+  }
+
+  SECTION("Kernel with arguments using kernelParams") {
+    LinearAllocGuard<int> result_dev(LinearAllocs::hipMalloc, sizeof(int));
+    HIP_CHECK(hipMemset(result_dev.ptr(), 0, sizeof(*result_dev.ptr())));
+    int* result_ptr = result_dev.ptr();
+    void* kernel_args[1] = {&result_ptr};
+    HIP_CHECK(hipExtLaunchKernel(reinterpret_cast<void*>(kernel_42), dim3{1, 1, 1}, dim3{1, 1, 1},
+                                 kernel_args, 0, nullptr, nullptr, nullptr, 0u));
+    int result = 0;
+    HIP_CHECK(hipMemcpy(&result, result_dev.ptr(), sizeof(result), hipMemcpyDefault));
+    REQUIRE(result == 42);
+  }
+}
+
+TEST_CASE("Unit_hipExtLaunchKernel_Positive_Parameters") {
+  SECTION("blockDim.x == maxBlockDimX") {
+    const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimX);
+    HIP_CHECK(hipExtLaunchKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1}, dim3{x, 1, 1},
+                                 nullptr, 0, nullptr, nullptr, nullptr, 0u));
+  }
+
+  SECTION("blockDim.y == maxBlockDimY") {
+    const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimY);
+    HIP_CHECK(hipExtLaunchKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1}, dim3{y, 1, 1},
+                                 nullptr, 0, nullptr, nullptr, nullptr, 0u));
+  }
+
+  SECTION("blockDim.z == maxBlockDimZ") {
+    const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimZ);
+    HIP_CHECK(hipExtLaunchKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1}, dim3{z, 1, 1},
+                                 nullptr, 0, nullptr, nullptr, nullptr, 0u));
+  }
+}
+
+TEST_CASE("Unit_hipExtLaunchKernel_Negative_Parameters") {
+  SECTION("f == nullptr") {
+    HIP_CHECK_ERROR(hipExtLaunchKernel(nullptr, dim3{1, 1, 1}, dim3{1, 1, 1}, nullptr, 0, nullptr,
+                                       nullptr, nullptr, 0u),
+                    hipErrorInvalidDeviceFunction);
+  }
+
+  SECTION("gridDim.x == 0") {
+    HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast<void*>(kernel), dim3{0, 1, 1},
+                                       dim3{1, 1, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("gridDim.y == 0") {
+    HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast<void*>(kernel), dim3{1, 0, 1},
+                                       dim3{1, 1, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("gridDim.z == 0") {
+    HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 0},
+                                       dim3{1, 1, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("blockDim.x == 0") {
+    HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
+                                       dim3{0, 1, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("blockDim.y == 0") {
+    HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
+                                       dim3{1, 0, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("blockDim.z == 0") {
+    HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
+                                       dim3{1, 1, 0}, nullptr, 0, nullptr, nullptr, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("blockDim.x > maxBlockDimX") {
+    const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimX) + 1u;
+    HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
+                                       dim3{x, 1, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u),
+                    hipErrorInvalidConfiguration);
+  }
+
+  SECTION("blockDim.y > maxBlockDimY") {
+    const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimY) + 1u;
+    HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
+                                       dim3{1, y, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u),
+                    hipErrorInvalidConfiguration);
+  }
+
+  SECTION("blockDim.z > maxBlockDimZ") {
+    const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimZ) + 1u;
+    HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
+                                       dim3{1, 1, z}, nullptr, 0, nullptr, nullptr, nullptr, 0u),
+                    hipErrorInvalidConfiguration);
+  }
+
+  SECTION("blockDim.x * blockDim.y * blockDim.z > maxThreadsPerBlock") {
+    const unsigned int max = GetDeviceAttribute(0, hipDeviceAttributeMaxThreadsPerBlock);
+    const unsigned int dim = std::ceil(std::cbrt(max));
+    HIP_CHECK_ERROR(
+        hipExtLaunchKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1}, dim3{dim, dim, dim},
+                           nullptr, 0, nullptr, nullptr, nullptr, 0u),
+        hipErrorInvalidConfiguration);
+  }
+
+  SECTION("sharedMemBytes > maxSharedMemoryPerBlock") {
+    const unsigned int max = GetDeviceAttribute(0, hipDeviceAttributeMaxSharedMemoryPerBlock) + 1u;
+    HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
+                                       dim3{1, 1, 1}, nullptr, max, nullptr, nullptr, nullptr, 0u),
+                    hipErrorOutOfMemory);
+  }
+
+  SECTION("Invalid stream") {
+    hipStream_t stream = nullptr;
+    HIP_CHECK(hipStreamCreate(&stream));
+    HIP_CHECK(hipStreamDestroy(stream));
+    HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
+                                       dim3{1, 1, 1}, nullptr, 0, stream, nullptr, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("Invalid startEvent") {
+    hipEvent_t event = nullptr;
+    HIP_CHECK(hipEventCreate(&event));
+    HIP_CHECK(hipEventDestroy(event));
+    HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
+                                       dim3{1, 1, 1}, nullptr, 0, nullptr, event, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("Invalid endEvent") {
+    hipEvent_t event = nullptr;
+    HIP_CHECK(hipEventCreate(&event));
+    HIP_CHECK(hipEventDestroy(event));
+    HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
+                                       dim3{1, 1, 1}, nullptr, 0, nullptr, nullptr, event, 0u),
+                    hipErrorInvalidValue);
+  }
+}
@@ -0,0 +1,144 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "execution_control_common.hh"
+
+#include <hip_test_common.hh>
+#include <hip/hip_runtime_api.h>
+#include <resource_guards.hh>
+#include <utils.hh>
+
+TEST_CASE("Unit_hipExtLaunchMultiKernelMultiDevice_Positive_Basic") {
+  const auto device_count = HipTest::getDeviceCount();
+
+  std::vector<hipLaunchParams> params_list(device_count);
+
+  int device = 0;
+  for (auto& params : params_list) {
+    params.func = reinterpret_cast<void*>(kernel);
+    params.gridDim = dim3{1, 1, 1};
+    params.blockDim = dim3{1, 1, 1};
+    params.args = nullptr;
+    params.sharedMem = 0;
+    HIP_CHECK(hipSetDevice(device++));
+    HIP_CHECK(hipStreamCreate(&params.stream));
+  }
+
+  HIP_CHECK(hipExtLaunchMultiKernelMultiDevice(params_list.data(), device_count, 0u));
+
+  for (const auto params : params_list) {
+    HIP_CHECK(hipStreamSynchronize(params.stream));
+  }
+
+  for (const auto params : params_list) {
+    HIP_CHECK(hipStreamDestroy(params.stream));
+  }
+}
+
+TEST_CASE("Unit_hipExtLaunchMultiKernelMultiDevice_Negative_Parameters") {
+  const auto device_count = HipTest::getDeviceCount();
+
+  std::vector<hipLaunchParams> params_list(device_count);
+
+  int device = 0;
+  for (auto& params : params_list) {
+    params.func = reinterpret_cast<void*>(kernel);
+    params.gridDim = dim3{1, 1, 1};
+    params.blockDim = dim3{1, 1, 1};
+    params.args = nullptr;
+    params.sharedMem = 0;
+    HIP_CHECK(hipSetDevice(device++));
+    HIP_CHECK(hipStreamCreate(&params.stream));
+  }
+
+  SECTION("launchParamsList == nullptr") {
+    HIP_CHECK_ERROR(hipExtLaunchMultiKernelMultiDevice(nullptr, device_count, 0u),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("numDevices == 0") {
+    HIP_CHECK_ERROR(hipExtLaunchMultiKernelMultiDevice(params_list.data(), 0, 0u),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("numDevices > device count") {
+    HIP_CHECK_ERROR(hipExtLaunchMultiKernelMultiDevice(params_list.data(), device_count + 1, 0u),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("invalid flags") {
+    HIP_CHECK_ERROR(hipExtLaunchMultiKernelMultiDevice(params_list.data(), device_count, 999),
+                    hipErrorInvalidValue);
+  }
+
+  if (device_count > 1) {
+    SECTION("launchParamsList.func doesn't match across all devices") {
+      params_list[1].func = reinterpret_cast<void*>(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));
+  }
+}
+
+TEST_CASE("Unit_hipExtLaunchMultiKernelMultiDevice_Negative_MultiKernelSameDevice") {
+  HIP_CHECK(hipSetDevice(0));
+
+  std::vector<hipLaunchParams> params_list(2);
+
+  for (auto& params : params_list) {
+    params.func = reinterpret_cast<void*>(kernel);
+    params.gridDim = dim3{1, 1, 1};
+    params.blockDim = dim3{1, 1, 1};
+    params.args = nullptr;
+    params.sharedMem = 0;
+    HIP_CHECK(hipStreamCreate(&params.stream));
+  }
+
+  HIP_CHECK_ERROR(hipExtLaunchMultiKernelMultiDevice(params_list.data(), 2, 0u),
+                  hipErrorInvalidValue);
+
+  for (const auto params : params_list) {
+    HIP_CHECK(hipStreamDestroy(params.stream));
+  }
+}