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 #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(HipTest::vectorADD), 0, 0)); + REQUIRE(gridSize > 0); + REQUIRE(blockSize > 0); + + HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor( + &numBlocks, reinterpret_cast(HipTest::vectorADD), + blockSize, 0)); + REQUIRE(numBlocks > 0); + + numBlocks = 0; + HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( + &numBlocks, reinterpret_cast(HipTest::vectorADD), + 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)); +} +