From 259010f2d5a3fca5bb65f22e43bf381837e84ac1 Mon Sep 17 00:00:00 2001 From: vstojilj Date: Thu, 27 Nov 2025 17:40:11 +0100 Subject: [PATCH] SWDEV-491253 - Create stream capture test for kernel APIs (#1189) --- .../executionControl/hipExtLaunchKernel.cc | 13 +++----- .../hipLaunchCooperativeKernel.cc | 19 +++++++++++ .../unit/executionControl/hipLaunchKernel.cc | 14 ++++++++ .../catch/unit/executionControl/launch_api.cc | 20 +++++++++++ .../hipModuleLaunchCooperativeKernel.cc | 33 +++++++++++++++++++ .../unit/module/hipModuleLaunchKernel.cc | 18 ++++++++++ 6 files changed, 108 insertions(+), 9 deletions(-) diff --git a/projects/hip-tests/catch/unit/executionControl/hipExtLaunchKernel.cc b/projects/hip-tests/catch/unit/executionControl/hipExtLaunchKernel.cc index be6ab767f7..b8f88da982 100644 --- a/projects/hip-tests/catch/unit/executionControl/hipExtLaunchKernel.cc +++ b/projects/hip-tests/catch/unit/executionControl/hipExtLaunchKernel.cc @@ -188,8 +188,6 @@ TEST_CASE("Unit_hipExtLaunchKernel_Negative_Parameters") { * - HIP_VERSION >= 6.0 */ TEST_CASE("Unit_hipExtLaunchKernel_capturehipExtLaunchKernel") { - hipGraph_t graph{nullptr}; - hipGraphExec_t graphExec{nullptr}; hipStream_t stream; HIP_CHECK(hipStreamCreate(&stream)); int* A_d; @@ -197,24 +195,21 @@ TEST_CASE("Unit_hipExtLaunchKernel_capturehipExtLaunchKernel") { A_h = reinterpret_cast(malloc(sizeof(int))); HIP_CHECK(hipMalloc(reinterpret_cast(&A_d), sizeof(int))); void* args[1] = {&A_d}; + // Begin Capture operation - HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + GENERATE_CAPTURE(); + BEGIN_CAPTURE(stream); HIP_CHECK(hipExtLaunchKernel(reinterpret_cast(kernel_42), dim3{1, 1, 1}, dim3{1, 1, 1}, args, 0, stream, nullptr, nullptr, 0u)); // End Capture - HIP_CHECK(hipStreamEndCapture(stream, &graph)); + END_CAPTURE(stream); - // Create and Launch Executable Graphs - HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); - HIP_CHECK(hipGraphLaunch(graphExec, stream)); HIP_CHECK(hipStreamSynchronize(stream)); HIP_CHECK(hipMemcpyDtoH(A_h, A_d, sizeof(int))); REQUIRE(A_h != nullptr); REQUIRE(*A_h == 42); - HIP_CHECK(hipGraphExecDestroy(graphExec)); - HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipStreamDestroy(stream)); HIP_CHECK(hipFree(A_d)); free(A_h); diff --git a/projects/hip-tests/catch/unit/executionControl/hipLaunchCooperativeKernel.cc b/projects/hip-tests/catch/unit/executionControl/hipLaunchCooperativeKernel.cc index 0684137866..a940c9fe8e 100644 --- a/projects/hip-tests/catch/unit/executionControl/hipLaunchCooperativeKernel.cc +++ b/projects/hip-tests/catch/unit/executionControl/hipLaunchCooperativeKernel.cc @@ -177,3 +177,22 @@ TEST_CASE("Unit_hipLaunchCooperativeKernel_Negative_Parameters") { hipErrorCooperativeLaunchTooLarge); } } + +TEST_CASE("Unit_hipLaunchCooperativeKernel_Verify_Capture") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + GENERATE_CAPTURE(); + BEGIN_CAPTURE(stream); + HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast(coop_kernel), dim3{2, 2, 1}, + dim3{1, 1, 1}, nullptr, 0, stream)); + END_CAPTURE(stream); + + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipStreamDestroy(stream)); +} diff --git a/projects/hip-tests/catch/unit/executionControl/hipLaunchKernel.cc b/projects/hip-tests/catch/unit/executionControl/hipLaunchKernel.cc index 030a12f5d9..c1ac883ba3 100644 --- a/projects/hip-tests/catch/unit/executionControl/hipLaunchKernel.cc +++ b/projects/hip-tests/catch/unit/executionControl/hipLaunchKernel.cc @@ -155,4 +155,18 @@ TEST_CASE("Unit_hipLaunchKernel_Negative_Parameters") { hipErrorInvalidValue); } #endif +} + +TEST_CASE("Unit_hipLaunchKernel_Verify_Capture") { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + GENERATE_CAPTURE(); + BEGIN_CAPTURE(stream); + HIP_CHECK(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, 1, 1}, nullptr, + 0, stream)); + END_CAPTURE(stream); + + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipStreamDestroy(stream)); } \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/executionControl/launch_api.cc b/projects/hip-tests/catch/unit/executionControl/launch_api.cc index 64cdcf8266..f301da7968 100644 --- a/projects/hip-tests/catch/unit/executionControl/launch_api.cc +++ b/projects/hip-tests/catch/unit/executionControl/launch_api.cc @@ -66,4 +66,24 @@ TEST_CASE("Unit___hipPushCallConfiguration_Positive_Basic") { REQUIRE(block.z == 1); REQUIRE(shmem == 1024); REQUIRE(stream == stream_guard.stream()); +} + +TEST_CASE("Unit_hipLaunchByPtr_Verify_Capture") { + LinearAllocGuard alloc(LinearAllocs::hipMallocManaged, 4); + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + GENERATE_CAPTURE(); + BEGIN_CAPTURE(stream); + + HIP_CHECK(hipConfigureCall(dim3{1}, dim3{1}, 0, stream)); + int* arg = alloc.ptr(); + HIP_CHECK(hipSetupArgument(&arg, sizeof(int*), 0)); + HIP_CHECK(hipLaunchByPtr(reinterpret_cast(kernel_42))); + + END_CAPTURE(stream); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipStreamDestroy(stream)); + + REQUIRE(alloc.ptr()[0] == 42); } \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/module/hipModuleLaunchCooperativeKernel.cc b/projects/hip-tests/catch/unit/module/hipModuleLaunchCooperativeKernel.cc index 6035989fc5..2c25b943af 100644 --- a/projects/hip-tests/catch/unit/module/hipModuleLaunchCooperativeKernel.cc +++ b/projects/hip-tests/catch/unit/module/hipModuleLaunchCooperativeKernel.cc @@ -213,6 +213,39 @@ TEST_CASE("Unit_hipModuleLaunchCooperativeKernel_Negative_Parameters") { #endif } +/** + * Test Description + * ------------------------ + * - Test `hipModuleLaunchCooperativeKernel` when it is captured. + * Test source + * ------------------------ + * - unit/module/hipModuleLaunchCooperativeKernel.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.5 + */ +TEST_CASE("Unit_hipModuleLaunchCooperativeKernel_Verify_Capture") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + + auto mg = ModuleGuard::InitModule("launch_kernel_module.code"); + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + GENERATE_CAPTURE(); + BEGIN_CAPTURE(stream); + + hipFunction_t f = GetKernel(mg.module(), "CoopKernel"); + HIP_CHECK(hipModuleLaunchCooperativeKernel(f, 2, 2, 1, 1, 1, 1, 0, stream, nullptr)); + + END_CAPTURE(stream); + + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipStreamDestroy(stream)); +} + /** * End doxygen group ModuleTest. * @} diff --git a/projects/hip-tests/catch/unit/module/hipModuleLaunchKernel.cc b/projects/hip-tests/catch/unit/module/hipModuleLaunchKernel.cc index 93c6a37962..f402de26e9 100644 --- a/projects/hip-tests/catch/unit/module/hipModuleLaunchKernel.cc +++ b/projects/hip-tests/catch/unit/module/hipModuleLaunchKernel.cc @@ -318,3 +318,21 @@ TEST_CASE("Unit_hipModuleLaunchKernel_Fntl") { REQUIRE(testStatus == true); } } + +TEST_CASE("Unit_hipModuleLaunchKernel_Verify_Capture") { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + auto mg = ModuleGuard::InitModule("launch_kernel_module.code"); + + GENERATE_CAPTURE(); + BEGIN_CAPTURE(stream); + + hipFunction_t f = GetKernel(mg.module(), "NOPKernel"); + HIP_CHECK(hipModuleLaunchKernel(f, 1, 1, 1, 1, 1, 1, 0, stream, nullptr, nullptr)); + + END_CAPTURE(stream); + + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipStreamDestroy(stream)); +}