EXSWHTEC-111 - Implement tests for the hipModuleLaunchKernel family of APIs #22

Change-Id: I963e17c413eb0976a1073e2f02a7e5eff1db3b42

catch/unit/module/CMakeLists.txt

@@ -26,13 +26,19 @@ set(TEST_SRC
     hipModuleLoadDataEx.cc
     hipModuleUnload.cc
     hipModuleGetFunction.cc
+    hipModuleLaunchKernel.cc
 )
 
 add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/get_function_module.code
                    COMMAND ${CMAKE_CXX_COMPILER} --genco --std=c++17 ${CMAKE_CURRENT_SOURCE_DIR}/get_function_module.cc -o get_function_module.code
-                   DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/get_function_module.cc)   
+                   DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/get_function_module.cc)
 add_custom_target(get_function_module ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/get_function_module.code)
 
+add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/launch_kernel_module.code
+                   COMMAND ${CMAKE_CXX_COMPILER} --genco --std=c++17 ${CMAKE_CURRENT_SOURCE_DIR}/launch_kernel_module.cc -o launch_kernel_module.code
+                   DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/launch_kernel_module.cc)
+add_custom_target(launch_kernel_module ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/launch_kernel_module.code)
+
 add_custom_target(empty_module.code
                   COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR}
                   ${CMAKE_CURRENT_SOURCE_DIR}/empty_module.cc
@@ -136,6 +142,7 @@ hip_add_exe_to_target(NAME ModuleTest
 
 add_dependencies(build_tests empty_module.code)
 add_dependencies(ModuleTest get_function_module)
+add_dependencies(ModuleTest launch_kernel_module)
 
 if(HIP_PLATFORM MATCHES "amd")
 add_dependencies(build_tests copyKernel.code copyKernel.s)
@@ -143,4 +150,4 @@ if(UNIX)
 add_dependencies(build_tests copiousArgKernel.code copiousArgKernel0.code copiousArgKernel1.code copiousArgKernel2.code
 copiousArgKernel3.code copiousArgKernel16.code copiousArgKernel17.code)
 endif()
-endif()
+endif()

catch/unit/module/hipExtModuleLaunchKernel.cc

@@ -50,6 +50,8 @@ THE SOFTWARE.
 #include "hip/hip_ext.h"
 #include <regex>  // NOLINT
 
+#include "hip_module_launch_kernel_common.hh"
+
 static constexpr auto totalWorkGroups{1024};
 static constexpr auto localWorkSize{512};
 static constexpr auto lastWorkSizeEven{256};
@@ -69,7 +71,7 @@ static bool searchRegExpr(const std::regex& expr, const char* filename) {
   assemblyfile.seekg(0, assemblyfile.end);
   int len = assemblyfile.tellg();
   assemblyfile.seekg(0, assemblyfile.beg);
-  char *fbuf = new char[len + 1];
+  char* fbuf = new char[len + 1];
   assemblyfile.read(fbuf, len);
   fbuf[len] = '\0';
 
@@ -124,8 +126,7 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_NonUniformWorkGroup") {
   auto isEven = GENERATE(0, 1);
   // Calculate size
   auto lastWorkSize = isEven ? lastWorkSizeEven : lastWorkSizeOdd;
-  size_t arraylength =
+  size_t arraylength = (totalWorkGroups - 1) * localWorkSize + lastWorkSize;
-  (totalWorkGroups - 1)*localWorkSize + lastWorkSize;
   size_t sizeBytes{arraylength * sizeof(int)};
   // Get module and function from module
   hipModule_t Module;
@@ -133,9 +134,9 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_NonUniformWorkGroup") {
   HIP_CHECK(hipModuleLoad(&Module, fileName));
   HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name));
   // Allocate resources
-  int *A = new int[arraylength];
+  int* A = new int[arraylength];
   REQUIRE(A != nullptr);
-  int *B = new int[arraylength];
+  int* B = new int[arraylength];
   REQUIRE(B != nullptr);
   // Inititialize data
   for (size_t i = 0; i < arraylength; i++) {
@@ -155,14 +156,13 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_NonUniformWorkGroup") {
   args.buffersize = arraylength;
   size_t size = sizeof(args);
 
-  void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
+  void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
-                    HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
                     HIP_LAUNCH_PARAM_END};
   // Memcpy from A to Ad
   HIP_CHECK(hipMemcpy(Ad, A, sizeBytes, hipMemcpyDefault));
-  REQUIRE(hipErrorInvalidValue == hipExtModuleLaunchKernel(Function,
+  REQUIRE(hipErrorInvalidValue ==
-                 arraylength, 1, 1, localWorkSize, 1, 1, 0, 0, NULL,
+          hipExtModuleLaunchKernel(Function, arraylength, 1, 1, localWorkSize, 1, 1, 0, 0, NULL,
-                 reinterpret_cast<void**>(&config), 0));
+                                   reinterpret_cast<void**>(&config), 0));
   HIP_CHECK(hipDeviceSynchronize());
   HIP_CHECK(hipFree(Ad));
   HIP_CHECK(hipFree(Bd));
@@ -194,9 +194,9 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_UniformWorkGroup") {
   HIP_CHECK(hipModuleLoad(&Module, fileName));
   HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name));
   // Allocate resources
-  int *A = new int[arraylength];
+  int* A = new int[arraylength];
   REQUIRE(A != nullptr);
-  int *B = new int[arraylength];
+  int* B = new int[arraylength];
   REQUIRE(B != nullptr);
   // Inititialize data
   for (size_t i = 0; i < arraylength; i++) {
@@ -216,14 +216,12 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_UniformWorkGroup") {
   args.buffersize = arraylength;
   size_t size = sizeof(args);
 
-  void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
+  void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
-                    HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
                     HIP_LAUNCH_PARAM_END};
   // Memcpy from A to Ad
   HIP_CHECK(hipMemcpy(Ad, A, sizeBytes, hipMemcpyDefault));
-  HIP_CHECK(hipExtModuleLaunchKernel(Function, arraylength, 1, 1,
+  HIP_CHECK(hipExtModuleLaunchKernel(Function, arraylength, 1, 1, localWorkSize, 1, 1, 0, 0, NULL,
-                                 localWorkSize, 1, 1, 0, 0, NULL,
+                                     reinterpret_cast<void**>(&config), 0));
-                                 reinterpret_cast<void**>(&config), 0));
   // Memcpy results back to host
   HIP_CHECK(hipMemcpy(B, Bd, sizeBytes, hipMemcpyDefault));
   HIP_CHECK(hipDeviceSynchronize());
@@ -237,3 +235,53 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_UniformWorkGroup") {
   delete[] B;
   HIP_CHECK(hipModuleUnload(Module));
 }
+
+TEST_CASE("Unit_hipExtModuleLaunchKernel_Positive_Basic") {
+  ModuleLaunchKernelPositiveBasic<hipExtModuleLaunchKernel>();
+
+  SECTION("Timed kernel launch with events") {
+    hipEvent_t start_event = nullptr, stop_event = nullptr;
+    HIP_CHECK(hipEventCreate(&start_event));
+    HIP_CHECK(hipEventCreate(&stop_event));
+    const auto kernel = GetKernel(mg.module(), "Delay");
+    int clock_rate = 0;
+    HIP_CHECK(hipDeviceGetAttribute(&clock_rate, hipDeviceAttributeClockRate, 0));
+    uint32_t interval = 100;
+    uint32_t ticks_per_second = clock_rate;
+    void* kernel_params[2] = {&interval, &ticks_per_second};
+    HIP_CHECK(hipExtModuleLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, nullptr, kernel_params, nullptr,
+                                       start_event, stop_event));
+    HIP_CHECK(hipDeviceSynchronize());
+    auto elapsed = 0.0f;
+    HIP_CHECK(hipEventElapsedTime(&elapsed, start_event, stop_event));
+    REQUIRE(static_cast<uint32_t>(elapsed) >= interval);
+  }
+}
+
+TEST_CASE("Unit_hipExtModuleLaunchKernel_Positive_Parameters") {
+  ModuleLaunchKernelPositiveParameters<hipExtModuleLaunchKernel>();
+
+  SECTION("Pass only start event") {
+    hipEvent_t start_event = nullptr;
+    HIP_CHECK(hipEventCreate(&start_event));
+    const auto kernel = GetKernel(mg.module(), "NOPKernel");
+    HIP_CHECK(hipExtModuleLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr,
+                                       start_event, nullptr));
+    HIP_CHECK(hipDeviceSynchronize());
+    HIP_CHECK(hipEventQuery(start_event));
+  }
+
+  SECTION("Pass only stop event") {
+    hipEvent_t stop_event = nullptr;
+    HIP_CHECK(hipEventCreate(&stop_event));
+    const auto kernel = GetKernel(mg.module(), "NOPKernel");
+    HIP_CHECK(hipExtModuleLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr,
+                                       nullptr, stop_event));
+    HIP_CHECK(hipDeviceSynchronize());
+    HIP_CHECK(hipEventQuery(stop_event));
+  }
+}
+
+TEST_CASE("Unit_hipExtModuleLaunchKernel_Negative_Parameters") {
+  ModuleLaunchKernelNegativeParameters<hipExtModuleLaunchKernel>();
+}

catch/unit/module/hipModuleLaunchKernel.cc

@@ -0,0 +1,49 @@
+/*
+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 "hip_module_launch_kernel_common.hh"
+
+#include <hip_test_common.hh>
+#include <hip/hip_runtime_api.h>
+
+static hipError_t hipModuleLaunchKernelWrapper(hipFunction_t f, uint32_t gridX, uint32_t gridY,
+                                               uint32_t gridZ, uint32_t blockX, uint32_t blockY,
+                                               uint32_t blockZ, size_t sharedMemBytes,
+                                               hipStream_t hStream, void** kernelParams,
+                                               void** extra, hipEvent_t, hipEvent_t, uint32_t) {
+  return hipModuleLaunchKernel(f, gridX, gridY, gridZ, blockX, blockY, blockZ, sharedMemBytes,
+                               hStream, kernelParams, extra);
+}
+
+TEST_CASE("Unit_hipModuleLaunchKernel_Positive_Basic") {
+  HIP_CHECK(hipFree(nullptr));
+  ModuleLaunchKernelPositiveBasic<hipModuleLaunchKernelWrapper>();
+}
+
+TEST_CASE("Unit_hipModuleLaunchKernel_Positive_Parameters") {
+  HIP_CHECK(hipFree(nullptr));
+  ModuleLaunchKernelPositiveParameters<hipModuleLaunchKernelWrapper>();
+}
+
+TEST_CASE("Unit_hipModuleLaunchKernel_Negative_Parameters") {
+  HIP_CHECK(hipFree(nullptr));
+  ModuleLaunchKernelNegativeParameters<hipModuleLaunchKernelWrapper>();
+}

catch/unit/module/hip_module_launch_kernel_common.hh

@@ -0,0 +1,269 @@
+/*
+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.
+*/
+
+#pragma once
+
+#include "hip_module_common.hh"
+
+#include <hip/hip_runtime_api.h>
+#include <resource_guards.hh>
+#include <utils.hh>
+
+inline int GetDeviceAttribute(const int device,
+                              const hipDeviceAttribute_t attr) {
+  int value = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&value, attr, device));
+  return value;
+}
+
+inline ModuleGuard InitModule() {
+  HIP_CHECK(hipFree(nullptr));
+  return ModuleGuard::LoadModule("launch_kernel_module.code");
+}
+
+inline ModuleGuard mg{InitModule()};
+
+using ExtModuleLaunchKernelSig = hipError_t(hipFunction_t, uint32_t, uint32_t, uint32_t, uint32_t,
+                                            uint32_t, uint32_t, size_t, hipStream_t, void**, void**,
+                                            hipEvent_t, hipEvent_t, uint32_t);
+
+template <ExtModuleLaunchKernelSig* func> void ModuleLaunchKernelPositiveBasic() {
+  SECTION("Kernel with no arguments") {
+    hipFunction_t f = GetKernel(mg.module(), "NOPKernel");
+    HIP_CHECK(func(f, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u));
+    HIP_CHECK(hipDeviceSynchronize());
+  }
+
+  SECTION("Kernel with arguments using kernelParams") {
+    hipFunction_t f = GetKernel(mg.module(), "Kernel42");
+    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(func(f, 1, 1, 1, 1, 1, 1, 0, nullptr, kernel_args, nullptr, nullptr, nullptr, 0u));
+    int result = 0;
+    HIP_CHECK(hipMemcpy(&result, result_dev.ptr(), sizeof(result), hipMemcpyDefault));
+    REQUIRE(result == 42);
+  }
+
+  SECTION("Kernel with arguments using extra") {
+    hipFunction_t f = GetKernel(mg.module(), "Kernel42");
+    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();
+    size_t size = sizeof(result_ptr);
+    // clang-format off
+    void *extra[] = {
+        HIP_LAUNCH_PARAM_BUFFER_POINTER, &result_ptr, 
+        HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, 
+        HIP_LAUNCH_PARAM_END
+    };
+    // clang-format on
+    HIP_CHECK(func(f, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr, extra, nullptr, nullptr, 0u));
+    int result = 0;
+    HIP_CHECK(hipMemcpy(&result, result_dev.ptr(), sizeof(result), hipMemcpyDefault));
+    REQUIRE(result == 42);
+  }
+}
+
+template <ExtModuleLaunchKernelSig* func> void ModuleLaunchKernelPositiveParameters() {
+  const auto LaunchNOPKernel = [=](unsigned int gridDimX, unsigned int gridDimY,
+                                   unsigned int gridDimZ, unsigned int blockDimX,
+                                   unsigned int blockDimY, unsigned int blockDimZ) {
+    hipFunction_t f = GetKernel(mg.module(), "NOPKernel");
+    HIP_CHECK(func(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, 0, nullptr,
+                   nullptr, nullptr, nullptr, nullptr, 0u));
+    HIP_CHECK(hipDeviceSynchronize());
+  };
+
+  SECTION("gridDimX == maxGridDimX") {
+    const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxGridDimX);
+    LaunchNOPKernel(x, 1, 1, 1, 1, 1);
+  }
+
+  SECTION("gridDimY == maxGridDimY") {
+    const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxGridDimY);
+    LaunchNOPKernel(1, y, 1, 1, 1, 1);
+  }
+
+  SECTION("gridDimZ == maxGridDimZ") {
+    const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxGridDimZ);
+    LaunchNOPKernel(1, 1, z, 1, 1, 1);
+  }
+
+  SECTION("blockDimX == maxBlockDimX") {
+    const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimX);
+    LaunchNOPKernel(1, 1, 1, x, 1, 1);
+  }
+
+  SECTION("blockDimY == maxBlockDimY") {
+    const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimY);
+    LaunchNOPKernel(1, 1, 1, 1, y, 1);
+  }
+
+  SECTION("blockDimZ == maxBlockDimZ") {
+    const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimZ);
+    LaunchNOPKernel(1, 1, 1, 1, 1, z);
+  }
+}
+
+template <ExtModuleLaunchKernelSig* func> void ModuleLaunchKernelNegativeParameters() {
+  hipFunction_t f = GetKernel(mg.module(), "NOPKernel");
+
+// Disabled on AMD due to defect - EXSWHTEC-157
+#if HT_NVIDIA
+  SECTION("f == nullptr") {
+    HIP_CHECK_ERROR(
+        func(nullptr, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u),
+        hipErrorInvalidResourceHandle);
+  }
+#endif
+
+  SECTION("gridDimX == 0") {
+    HIP_CHECK_ERROR(func(f, 0, 1, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("gridDimY == 0") {
+    HIP_CHECK_ERROR(func(f, 1, 0, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("gridDimZ == 0") {
+    HIP_CHECK_ERROR(func(f, 1, 1, 0, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("blockDimX == 0") {
+    HIP_CHECK_ERROR(func(f, 1, 1, 1, 0, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("blockDimY == 0") {
+    HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, 0, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("blockDimZ == 0") {
+    HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, 1, 0, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+
+// Disabled on AMD due to defect - EXSWHTEC-158
+#if HT_NVIDIA
+  SECTION("gridDimX > maxGridDimX") {
+    const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxGridDimX) + 1u;
+    HIP_CHECK_ERROR(func(f, x, 1, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("gridDimY > maxGridDimY") {
+    const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxGridDimY) + 1u;
+    HIP_CHECK_ERROR(func(f, 1, y, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("gridDimZ > maxGridDimZ") {
+    const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxGridDimZ) + 1u;
+    HIP_CHECK_ERROR(func(f, 1, 1, z, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+#endif
+
+// Disabled on AMD due to defect - EXSWHTEC-156
+#if HT_NVIDIA
+  SECTION("blockDimX > maxBlockDimX") {
+    const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimX) + 1u;
+    HIP_CHECK_ERROR(func(f, 1, 1, 1, x, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("blockDimY > maxBlockDimY") {
+    const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimY) + 1u;
+    HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, y, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("blockDimZ > maxBlockDimZ") {
+    const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimZ) + 1u;
+    HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, 1, z, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+#endif
+
+// Disabled on AMD due to defect - EXSWHTEC-162
+#if HT_NVIDIA
+  SECTION("blockDimX * blockDimY * blockDimZ > MaxThreadsPerBlock") {
+    const unsigned int max = GetDeviceAttribute(0, hipDeviceAttributeMaxThreadsPerBlock);
+    const unsigned int dim = std::ceil(std::cbrt(max)) + 1;
+    HIP_CHECK_ERROR(
+        func(f, 1, 1, 1, dim, dim, dim, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u),
+        hipErrorInvalidValue);
+  }
+#endif
+
+// Disabled on AMD due to defect - EXSWHTEC-159
+#if HT_NVIDIA
+  SECTION("sharedMemBytes > max shared memory per block") {
+    const unsigned int max = GetDeviceAttribute(0, hipDeviceAttributeMaxSharedMemoryPerBlock) + 1u;
+    HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, 1, 1, max, nullptr, nullptr, nullptr, nullptr, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+#endif
+
+// Disabled on AMD due to defect - EXSWHTEC-160
+#if HT_NVIDIA
+  SECTION("Invalid stream") {
+    hipStream_t stream = nullptr;
+    HIP_CHECK(hipStreamCreate(&stream));
+    HIP_CHECK(hipStreamDestroy(stream));
+    HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, 1, 0, 0, stream, nullptr, nullptr, nullptr, nullptr, 0u),
+                    hipErrorContextIsDestroyed);
+  }
+#endif
+
+  SECTION("Passing kernel_args and extra simultaneously") {
+    hipFunction_t f = GetKernel(mg.module(), "Kernel42");
+    LinearAllocGuard<int> result_dev(LinearAllocs::hipMalloc, sizeof(int));
+    int* result_ptr = result_dev.ptr();
+    size_t size = sizeof(result_ptr);
+    void* kernel_args[1] = {&result_ptr};
+    // clang-format off
+    void *extra[] = {
+        HIP_LAUNCH_PARAM_BUFFER_POINTER, &result_ptr, 
+        HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, 
+        HIP_LAUNCH_PARAM_END
+    };
+    // clang-format on
+    HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, 1, 1, 0, nullptr, kernel_args, extra, nullptr, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+
+// Disabled on AMD due to defect - EXSWHTEC-161
+#if HT_NVIDIA
+  SECTION("Invalid extra") {
+    hipFunction_t f = GetKernel(mg.module(), "Kernel42");
+    void* extra[0] = {};
+    HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr, extra, nullptr, nullptr, 0u),
+                    hipErrorInvalidValue);
+  }
+#endif
+}

catch/unit/module/launch_kernel_module.cc

@@ -0,0 +1,37 @@
+/*
+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 <hip/hip_runtime.h>
+
+extern "C" {
+__global__ void NOPKernel() {}
+
+__global__ void Kernel42(int* out) { *out = 42; }
+
+// Interval is in millisecond
+__global__ void Delay(uint32_t interval, const uint32_t ticks_per_ms) {
+  while (interval--) {
+    uint64_t start = clock();
+    while (clock() - start < ticks_per_ms) {
+    }
+  }
+}
+}