SWDEV-508349-[catch2][dtest] _spt api names missing from code coverage(phase-2)

catch/unit/event/CMakeLists.txt

@@ -10,6 +10,7 @@ set(TEST_SRC
     hipEventCreateWithFlags.cc
     hipEventSynchronize.cc
     Unit_hipEventMGpuMThreads.cc
+    hipEventRecord_spt.cc
 )
 
 # The test used wait mechanism and doesnt play well with all arch of nvidia

catch/unit/event/hipEventRecord_spt.cc

@@ -0,0 +1,151 @@
+/*Copyright (c) 2025 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_test_common.hh>
+#include <kernels.hh>
+#include <hip_test_checkers.hh>
+#include <hip_test_context.hh>
+#include <hip_test_defgroups.hh>
+/**
+ * @addtogroup hipEventRecord_spt hipEventRecord_spt
+ * @{
+ * @ingroup EventTest
+ * `hipEventRecord_spt(hipEvent_t event, hipStream_t stream = NULL)` -
+ * Record an event in the specified stream.
+ */
+/**
+ * Test Description
+ * ------------------------
+ *  - Creates regular events and events with flags.
+ *  - Enqueues them to the streams and checks if events
+ *    can be successfully used for synchronization.
+ * Test source
+ * ------------------------
+ *  - unit/event/hipEventRecord_spt.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipEventRecord_spt_BasicTst") {
+  constexpr size_t N = 1024;
+  constexpr int iterations = 1;
+  constexpr int blocks = 1024;
+  constexpr size_t Nbytes = N * sizeof(float);
+  float *A_h, *B_h, *C_h;
+  float *A_d, *B_d, *C_d;
+  HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N);
+  enum TestType {
+    WithFlags_Default = hipEventDefault,
+    WithFlags_Blocking = hipEventBlockingSync,
+    WithFlags_DisableTiming = hipEventDisableTiming,
+#if HT_AMD
+    WithFlags_ReleaseToDevice = hipEventReleaseToDevice,
+    WithFlags_ReleaseToSystem = hipEventReleaseToSystem,
+#endif
+    WithoutFlags
+  };
+#if HT_AMD
+  auto flags = GENERATE(WithFlags_Default, WithFlags_Blocking, WithFlags_DisableTiming,
+                        WithFlags_ReleaseToDevice, WithFlags_ReleaseToSystem, WithoutFlags);
+#endif
+#if HT_NVIDIA
+  auto flags =
+      GENERATE(WithFlags_Default, WithFlags_Blocking, WithFlags_DisableTiming, WithoutFlags);
+#endif
+  hipEvent_t start{}, stop{};
+  if (flags == WithoutFlags) {
+    HIP_CHECK(hipEventCreate(&start));
+    HIP_CHECK(hipEventCreate(&stop));
+  } else {
+    HIP_CHECK(hipEventCreateWithFlags(&start, flags));
+    HIP_CHECK(hipEventCreateWithFlags(&stop, flags));
+  }
+  HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
+  // Warmup
+  HipTest::launchKernel<float>(HipTest::vectorADD<float>, blocks, 1, 0, 0,
+                               static_cast<const float*>(A_d), static_cast<const float*>(B_d), C_d,
+                               N);
+  HIP_CHECK(hipDeviceSynchronize());
+  for (int i = 0; i < iterations; i++) {
+    //--- START TIMED REGION
+    long long hostStart = HipTest::get_time();
+    // Record the start event
+    HIP_CHECK(hipEventRecord_spt(start, NULL));
+    HipTest::launchKernel<float>(HipTest::vectorADD<float>, blocks, 1, 0, 0,
+                                 static_cast<const float*>(A_d), static_cast<const float*>(B_d),
+                                 C_d, N);
+    HIP_CHECK(hipGetLastError());
+    HIP_CHECK(hipEventRecord_spt(stop, NULL));
+    HIP_CHECK(hipEventSynchronize(stop));
+    long long hostStop = HipTest::get_time();
+    //--- STOP TIMED REGION
+    float hostMs = HipTest::elapsed_time(hostStart, hostStop);
+    INFO("host_time (chrono)                = " << hostMs);
+    // Make sure timer is timing something...
+    if (flags != WithFlags_DisableTiming) {
+      float eventMs = 1.0f;
+      HIP_CHECK(hipEventElapsedTime(&eventMs, start, stop));
+      INFO("kernel_time (hipEventElapsedTime) = " << eventMs);
+      REQUIRE(eventMs > 0.0f);
+    }
+  }
+  HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
+  HIP_CHECK(hipEventDestroy(start));
+  HIP_CHECK(hipEventDestroy(stop));
+  HipTest::checkVectorADD(A_h, B_h, C_h, N, true);
+  HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
+  TestContext::get().cleanContext();
+}
+/**
+ * Test Description
+ * ------------------------
+ *  - Validates handling of invalid arguments:
+ *    -# When event is `nullptr`
+ *      - Expected output: return `hipErrorInvalidResourceHandle`
+ *    -# When event is created on one device but recorded on the other one
+ *      - Expected output: return `hipErrorInvalidHandle`
+ * Test source
+ * ------------------------
+ *  - unit/event/hipEventRecord_spt.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipEventRecord_spt_Negative") {
+  SECTION("Nullptr event") {
+    HIP_CHECK_ERROR(hipEventRecord_spt(nullptr, nullptr), hipErrorInvalidResourceHandle);
+  }
+  SECTION("Different devices") {
+    int devCount = 0;
+    HIP_CHECK(hipGetDeviceCount(&devCount));
+    if (devCount > 1) {
+      // create event on dev=0
+      HIP_CHECK(hipSetDevice(0));
+      hipEvent_t start;
+      HIP_CHECK(hipEventCreate(&start));
+      // start on device 0 but null stream on device 1
+      HIP_CHECK(hipSetDevice(1));
+      HIP_CHECK_ERROR(hipEventRecord_spt(start, nullptr), hipErrorInvalidHandle);
+      HIP_CHECK(hipEventDestroy(start));
+    }
+  }
+}
+/**
+ * End doxygen group EventTest.
+ * @}
+ */

catch/unit/executionControl/CMakeLists.txt

@@ -12,6 +12,7 @@ if(HIP_PLATFORM MATCHES "amd")
         hipExtLaunchMultiKernelMultiDevice.cc
         launch_api.cc
         hipGetProcAddressLaunchCbExecCtrlApis.cc
+	hipLaunchCooperativeKernel_spt.cc
     )
 else()
     # These functions are currently unimplemented on AMD

catch/unit/executionControl/hipLaunchCooperativeKernel_spt.cc

@@ -0,0 +1,201 @@
+/*Copyright (c) 2025 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_test_common.hh>
+#include <hip/hip_runtime_api.h>
+#include <resource_guards.hh>
+#include <hip_test_defgroups.hh>
+#include <utils.hh>
+#include "execution_control_common.hh"
+/**
+ * @addtogroup hipLaunchCooperativeKernel_spt hipLaunchCooperativeKernel_spt
+ * @{
+ * @ingroup ExecutionTest
+ * `hipError_t hipLaunchCooperativeKernel_spt(const void* f, dim3 gridDim, dim3 blockDimX,
+                                      void** kernelParams, unsigned int sharedMemBytes,
+                                      hipStream_t stream);` -
+ * launches kernel f with launch parameters and shared memory on stream with arguments passed
+ * to kernelparams or extra, where thread blocks can cooperate and synchronize as they execute.
+ */
+/**
+ * Test Description
+ * ------------------------
+ * - Basic test to check the functionality of hipLaunchCooperativeKernel_spt.
+ * Test source
+ * ------------------------
+ * - catch\unit\executionControl\hipLaunchCooperativeKernel_spt.cc
+ * Test requirements
+ * ------------------------
+ * - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipLaunchCooperativeKernel_spt_Positive_Basic") {
+  if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) {
+    HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported");
+    return;
+  }
+  SECTION("Cooperative kernel with no arguments") {
+    HIP_CHECK(hipLaunchCooperativeKernel_spt(reinterpret_cast<void*>(coop_kernel), dim3{2, 2, 1},
+                                             dim3{1, 1, 1}, nullptr, 0, nullptr));
+    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(hipLaunchCooperativeKernel_spt(reinterpret_cast<void*>(kernel_42), dim3{1, 1, 1},
+                                             dim3{1, 1, 1}, kernel_args, 0, nullptr));
+    int result = 0;
+    HIP_CHECK(hipMemcpy(&result, result_dev.ptr(), sizeof(result), hipMemcpyDefault));
+    REQUIRE(result == 42);
+  }
+}
+/**
+ * Test Description
+ * ------------------------
+ * - Basic test to check the functionality of hipLaunchCooperativeKernel_spt
+ * with positive parameters.
+ * Test source
+ * ------------------------
+ * - catch\unit\executionControl\hipLaunchCooperativeKernel_spt.cc
+ * Test requirements
+ * ------------------------
+ * - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipLaunchCooperativeKernel_spt_Positive_Parameters") {
+  if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) {
+    HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported");
+    return;
+  }
+  SECTION("blockDim.x == maxBlockDimX") {
+    const unsigned int x = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimX, 0);
+    HIP_CHECK(hipLaunchCooperativeKernel_spt(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
+                                             dim3{x, 1, 1}, nullptr, 0, nullptr));
+  }
+  SECTION("blockDim.y == maxBlockDimY") {
+    const unsigned int y = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimY, 0);
+    HIP_CHECK(hipLaunchCooperativeKernel_spt(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
+                                             dim3{y, 1, 1}, nullptr, 0, nullptr));
+  }
+  SECTION("blockDim.z == maxBlockDimZ") {
+    const unsigned int z = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimZ, 0);
+    HIP_CHECK(hipLaunchCooperativeKernel_spt(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
+                                             dim3{z, 1, 1}, nullptr, 0, nullptr));
+  }
+}
+/**
+ * Test Description
+ * ------------------------
+ * - Basic test to check the functionality of hipLaunchCooperativeKernel_spt
+ * with negative parameters.
+ * Test source
+ * ------------------------
+ * - catch\unit\executionControl\hipLaunchCooperativeKernel_spt.cc
+ * Test requirements
+ * ------------------------
+ * - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipLaunchCooperativeKernel_spt_Negative_Parameters") {
+  if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) {
+    HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported");
+    return;
+  }
+  SECTION("f == nullptr") {
+    HIP_CHECK_ERROR(hipLaunchCooperativeKernel_spt(static_cast<void*>(nullptr), dim3{1, 1, 1},
+                                                   dim3{1, 1, 1}, nullptr, 0, nullptr),
+                    hipErrorInvalidDeviceFunction);
+  }
+  SECTION("gridDim.x == 0") {
+    HIP_CHECK_ERROR(hipLaunchCooperativeKernel_spt(reinterpret_cast<void*>(kernel), dim3{0, 1, 1},
+                                                   dim3{1, 1, 1}, nullptr, 0, nullptr),
+                    hipErrorInvalidConfiguration);
+  }
+  SECTION("gridDim.y == 0") {
+    HIP_CHECK_ERROR(hipLaunchCooperativeKernel_spt(reinterpret_cast<void*>(kernel), dim3{1, 0, 1},
+                                                   dim3{1, 1, 1}, nullptr, 0, nullptr),
+                    hipErrorInvalidConfiguration);
+  }
+  SECTION("gridDim.z == 0") {
+    HIP_CHECK_ERROR(hipLaunchCooperativeKernel_spt(reinterpret_cast<void*>(kernel), dim3{1, 1, 0},
+                                                   dim3{1, 1, 1}, nullptr, 0, nullptr),
+                    hipErrorInvalidConfiguration);
+  }
+  SECTION("blockDim.x == 0") {
+    HIP_CHECK_ERROR(hipLaunchCooperativeKernel_spt(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
+                                                   dim3{0, 1, 1}, nullptr, 0, nullptr),
+                    hipErrorInvalidConfiguration);
+  }
+  SECTION("blockDim.y == 0") {
+    HIP_CHECK_ERROR(hipLaunchCooperativeKernel_spt(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
+                                                   dim3{1, 0, 1}, nullptr, 0, nullptr),
+                    hipErrorInvalidConfiguration);
+  }
+  SECTION("blockDim.z == 0") {
+    HIP_CHECK_ERROR(hipLaunchCooperativeKernel_spt(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
+                                                   dim3{1, 1, 0}, nullptr, 0, nullptr),
+                    hipErrorInvalidConfiguration);
+  }
+  SECTION("blockDim.x > maxBlockDimX") {
+    const unsigned int x = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimX, 0) + 1u;
+    HIP_CHECK_ERROR(hipLaunchCooperativeKernel_spt(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
+                                                   dim3{x, 1, 1}, nullptr, 0, nullptr),
+                    hipErrorInvalidConfiguration);
+  }
+  SECTION("blockDim.y > maxBlockDimY") {
+    const unsigned int y = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimY, 0) + 1u;
+    HIP_CHECK_ERROR(hipLaunchCooperativeKernel_spt(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
+                                                   dim3{1, y, 1}, nullptr, 0, nullptr),
+                    hipErrorInvalidConfiguration);
+  }
+  SECTION("blockDim.z > maxBlockDimZ") {
+    const unsigned int z = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimZ, 0) + 1u;
+    HIP_CHECK_ERROR(hipLaunchCooperativeKernel_spt(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
+                                                   dim3{1, 1, z}, nullptr, 0, nullptr),
+                    hipErrorInvalidConfiguration);
+  }
+  SECTION("blockDim.x * blockDim.y * blockDim.z > maxThreadsPerBlock") {
+    const unsigned int max = GetDeviceAttribute(hipDeviceAttributeMaxThreadsPerBlock, 0);
+    const unsigned int dim = std::ceil(std::cbrt(max));
+    HIP_CHECK_ERROR(hipLaunchCooperativeKernel_spt(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
+                                                   dim3{dim, dim, dim}, nullptr, 0, nullptr),
+                    hipErrorInvalidConfiguration);
+  }
+  SECTION(
+      "gridDim.x * gridDim.y * gridDim.z > maxActiveBlocksPerMultiprocessor * "
+      "multiProcessorCount") {
+    int max_blocks;
+    HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks,
+                                                           reinterpret_cast<void*>(kernel), 1, 0));
+    const unsigned int multiproc_count =
+        GetDeviceAttribute(hipDeviceAttributeMultiprocessorCount, 0);
+    const unsigned int dim = std::ceil(std::cbrt(max_blocks * multiproc_count));
+    HIP_CHECK_ERROR(
+        hipLaunchCooperativeKernel_spt(reinterpret_cast<void*>(kernel), dim3{dim, dim, dim},
+                                       dim3{1, 1, 1}, nullptr, 0, nullptr),
+        hipErrorCooperativeLaunchTooLarge);
+  }
+  SECTION("sharedMemBytes > maxSharedMemoryPerBlock") {
+    const unsigned int max = GetDeviceAttribute(hipDeviceAttributeMaxSharedMemoryPerBlock, 0) + 1u;
+    HIP_CHECK_ERROR(hipLaunchCooperativeKernel_spt(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
+                                                   dim3{1, 1, 1}, nullptr, max, nullptr),
+                    hipErrorCooperativeLaunchTooLarge);
+  }
+}
+/**
+ * End doxygen group ExecutionTest.
+ * @}
+ */

catch/unit/graph/CMakeLists.txt

@@ -93,6 +93,12 @@ if(HIP_PLATFORM MATCHES "amd")
     hipStreamCaptureExtModuleLaunchKernel.cc
     hipStreamBeginCaptureToGraph.cc
     hipGetProcAddressGraphApis.cc
+    hipStreamGetCaptureInfo_spt.cc
+    hipStreamGetCaptureInfo_v2_spt.cc
+    hipStreamIsCapturing_spt.cc
+    hipStreamBeginCapture_spt.cc
+    hipStreamEndCapture_spt.cc
+    hipGraphLaunch_spt.cc
     # Below files are disbled in NVIDIA as PSDB builds are failing due to lower CUDA version.
     hipGraphExecNodeSetParams.cc
     hipGraphNodeSetParams.cc