From ea366c7aa5a8bd2b24e13a12bc06b364740d63eb Mon Sep 17 00:00:00 2001
From: "Swargam, Rambabu" <Rambabu.Swargam@amd.com>
Date: Mon, 7 Jul 2025 17:42:45 +0530
Subject: [PATCH] SWDEV-491288 - [catch2][dtest] Stream Capture tests for
 Occupancy API's (#157)

* SWDEV-491288 - [catch2][dtest] Stream Capture tests for Occupancy API's

* Modified Cmake file
---
 catch/unit/occupancy/CMakeLists.txt           |  2 +-
 ...cupancyMaxActiveBlocksPerMultiprocessor.cc | 72 +++++++++++++++++++
 2 files changed, 73 insertions(+), 1 deletion(-)

diff --git a/catch/unit/occupancy/CMakeLists.txt b/catch/unit/occupancy/CMakeLists.txt
index e7cdd4d6da..3f5a3b9827 100644
--- a/catch/unit/occupancy/CMakeLists.txt
+++ b/catch/unit/occupancy/CMakeLists.txt
@@ -13,7 +13,7 @@ set(TEST_SRC
 )
 
 add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/simple_kernel.code
-                   COMMAND ${CMAKE_CXX_COMPILER} --genco --std=c++17
+                   COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} --std=c++17
                    ${CMAKE_CURRENT_SOURCE_DIR}/simple_kernel.cc
                    -I${HIP_PATH}/include/
                    -o simple_kernel.code --hip-path=${HIP_PATH}
diff --git a/catch/unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessor.cc b/catch/unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessor.cc
index 65e1b42ab3..4dde9620a7 100644
--- a/catch/unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessor.cc
+++ b/catch/unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessor.cc
@@ -23,6 +23,7 @@ execution of hipModuleOccupancyMaxActiveBlocksPerMultiprocessor for diffrent par
 Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters - Test unsuccessful
 execution of hipModuleOccupancyMaxActiveBlocksPerMultiprocessor api when parameters are invalid
 */
+#include <hip_test_kernels.hh>
 #include "occupancy_common.hh"
 
 TEST_CASE("Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters") {
@@ -90,3 +91,74 @@ TEST_CASE("Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_Positive_Rang
 
   HIP_CHECK(hipModuleUnload(module));
 }
+
+/**
+ * Test Description
+ * ------------------------
+ *  - This test case tests the behaviour of all Occupancy APIs
+ *  - during the the stream capture.
+ * Test source
+ * ------------------------
+ *  - unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessor.cc
+ */
+TEST_CASE("Unit_OccupancyAPIs_StreamCapture") {
+  GENERATE_CAPTURE();
+
+  hipStream_t stream;
+  HIP_CHECK(hipStreamCreate(&stream));
+
+  hipModule_t module;
+  HIP_CHECK(hipModuleLoad(&module, "simple_kernel.code"));
+  REQUIRE(module != nullptr);
+  hipFunction_t function;
+  HIP_CHECK(hipModuleGetFunction(&function, module, "SimpleKernel"));
+  REQUIRE(function != nullptr);
+
+  BEGIN_CAPTURE(stream);
+
+  int gridSize = 0, blockSize = 0, numBlocks = 0;
+
+  HIP_CHECK(hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize,
+            reinterpret_cast<const void *>(HipTest::vectorADD<int>), 0, 0));
+  REQUIRE(gridSize > 0);
+  REQUIRE(blockSize > 0);
+
+  HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(
+            &numBlocks, reinterpret_cast<const void *>(HipTest::vectorADD<int>),
+            blockSize, 0));
+  REQUIRE(numBlocks > 0);
+
+  numBlocks = 0;
+  HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
+            &numBlocks, reinterpret_cast<const void *>(HipTest::vectorADD<int>),
+            blockSize, 0, 0));
+  REQUIRE(numBlocks > 0);
+
+  gridSize = 0, blockSize = 0;
+  HIP_CHECK(hipModuleOccupancyMaxPotentialBlockSize(&gridSize, &blockSize,
+                                                    function, 0, 0));
+  REQUIRE(gridSize > 0);
+  REQUIRE(blockSize > 0);
+
+  gridSize = 0; blockSize = 0;
+  HIP_CHECK(hipModuleOccupancyMaxPotentialBlockSizeWithFlags(
+            &gridSize, &blockSize, function, 0, 0, 0));
+  REQUIRE(gridSize > 0);
+  REQUIRE(blockSize > 0);
+
+  numBlocks = 0;
+  HIP_CHECK(hipModuleOccupancyMaxActiveBlocksPerMultiprocessor(
+            &numBlocks, function, blockSize, 0));
+  REQUIRE(numBlocks > 0);
+
+  numBlocks = 0;
+  HIP_CHECK(hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
+            &numBlocks, function, blockSize, 0, 0));
+  REQUIRE(numBlocks > 0);
+
+  END_CAPTURE(stream);
+
+  HIP_CHECK(hipModuleUnload(module));
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+