diff --git a/projects/hip/tests/catch/unit/stream/CMakeLists.txt b/projects/hip/tests/catch/unit/stream/CMakeLists.txt index 14c05e7f74..31157f7da6 100644 --- a/projects/hip/tests/catch/unit/stream/CMakeLists.txt +++ b/projects/hip/tests/catch/unit/stream/CMakeLists.txt @@ -9,6 +9,8 @@ set(TEST_SRC hipStreamCreateWithFlags.cc hipStreamCreateWithPriority.cc hipStreamWithCUMask.cc + hipStreamGetCUMask.cc + hipAPIStreamDisable.cc ) else() set(TEST_SRC @@ -20,6 +22,7 @@ set(TEST_SRC hipStreamAddCallback.cc hipStreamCreateWithFlags.cc hipStreamCreateWithPriority.cc + hipAPIStreamDisable.cc ) endif() diff --git a/projects/hip/tests/catch/unit/stream/hipAPIStreamDisable.cc b/projects/hip/tests/catch/unit/stream/hipAPIStreamDisable.cc new file mode 100644 index 0000000000..4da259074b --- /dev/null +++ b/projects/hip/tests/catch/unit/stream/hipAPIStreamDisable.cc @@ -0,0 +1,68 @@ +/* +Copyright (c) 2021 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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include +#include "hip/math_functions.h" + +#define NUM_STREAMS 8 + +namespace hipAPIStreamDisableTest { +const int NN = 1 << 21; + +__global__ void kernel(float* x, float* y, int n) { + int tid = threadIdx.x; + if (tid < 1) { + for (int i = 0; i < n; i++) { + x[i] = sqrt(powf(3.14159, i)); + } + y[tid] = y[tid] + 1.0f; + } +} + +__global__ void nKernel(float* y) { + int tid = threadIdx.x; + y[tid] = y[tid] + 1.0f; +} +} // namespace hipAPIStreamDisableTest + +/** + * Validate basic multistream functionalities + */ +TEST_CASE("Unit_hipStreamCreate_MultistreamBasicFunctionalities") { + hipStream_t streams[NUM_STREAMS]; + float *data[NUM_STREAMS], *yd, *xd; + float y = 1.0f, x = 1.0f; + HIP_CHECK(hipMalloc(reinterpret_cast(&yd), sizeof(float))); + HIP_CHECK(hipMalloc(reinterpret_cast(&xd), sizeof(float))); + HIP_CHECK(hipMemcpy(yd, &y, sizeof(float), hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(xd, &x, sizeof(float), hipMemcpyHostToDevice)); + + for (int i = 0; i < NUM_STREAMS; i++) { + HIP_CHECK(hipStreamCreate(&streams[i])); + HIP_CHECK(hipMalloc(&data[i], + (hipAPIStreamDisableTest::NN)*sizeof(float))); + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipAPIStreamDisableTest::kernel), + dim3(1), dim3(1), 0, streams[i], data[i], xd, + hipAPIStreamDisableTest::NN); + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipAPIStreamDisableTest::nKernel), + dim3(1), dim3(1), 0, 0, yd); + } + HIP_CHECK(hipMemcpy(&x, xd, sizeof(float), hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(&y, yd, sizeof(float), hipMemcpyDeviceToHost)); + REQUIRE(x == y); +} diff --git a/projects/hip/tests/catch/unit/stream/hipStreamGetCUMask.cc b/projects/hip/tests/catch/unit/stream/hipStreamGetCUMask.cc new file mode 100644 index 0000000000..49f21bdbda --- /dev/null +++ b/projects/hip/tests/catch/unit/stream/hipStreamGetCUMask.cc @@ -0,0 +1,177 @@ +/* +Copyright (c) 2021 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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Testcase Scenarios : +1) Test to verify hipExtStreamGetCUMask api returning default CU Mask or global CU Mask. +2) Test to verify hipExtStreamGetCUMask api returns custom mask set. +3) Negative tests for hipExtStreamGetCUMask api. +*/ + +#include +#include + + +/** + * Scenario to verify hipExtStreamGetCUMask api returning default CU Mask or global CU Mask. + * Scenario to verify hipExtStreamGetCUMask api returns custom mask set. + */ +TEST_CASE("Unit_hipExtStreamGetCUMask_verifyDefaultAndCustomMask") { + constexpr int maxNum = 6; + std::vector cuMask(maxNum); + hipDeviceProp_t props; + std::stringstream ss; + char* gCUMask{nullptr}; + std::string globalCUMask(""); + std::vector defaultCUMask; + + int nGpu = 0; + HIP_CHECK(hipGetDeviceCount(&nGpu)); + if (nGpu < 1) { + INFO("info: didn't find any GPU! skipping the test!"); + return; + } + + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipGetDeviceProperties(&props, 0)); + INFO("info: running on bus " << "0x" << props.pciBusID << " " << + props.name << " with " << props.multiProcessorCount << " CUs"); + + // Get global CU Mask if exists + gCUMask = getenv("ROC_GLOBAL_CU_MASK"); + if (gCUMask != nullptr && gCUMask[0] != '\0') { + globalCUMask.assign(gCUMask); + + for_each(globalCUMask.begin(), globalCUMask.end(), [](char & c) { + c = ::tolower(c); + }); + } + + // Create default CU Mask + uint32_t temp = 0; + uint32_t bit_index = 0; + for (uint32_t i = 0; i < (uint32_t)props.multiProcessorCount; i++) { + temp |= 1UL << bit_index; + if (bit_index >= 32) { + defaultCUMask.push_back(temp); + temp = 0; + bit_index = 0; + temp |= 1UL << bit_index; + } + bit_index += 1; + } + if (bit_index != 0) { + defaultCUMask.push_back(temp); + } + + SECTION("Verify with default CU Mask or global CU Mask") { + // make a default CU mask bit-array where all CUs are active + // this default mask is expected to be returned when there is no + // custom or global CU mask defined + + HIP_CHECK(hipExtStreamGetCUMask(0, cuMask.size(), &cuMask[0])); + + ss << std::hex; + for (int i = cuMask.size() - 1; i >= 0; i--) { + ss << cuMask[i]; + } + + // remove extra 0 from ss if any present + size_t found = ss.str().find_first_not_of("0"); + if (found != std::string::npos) { + ss.str(ss.str().substr(found, ss.str().length())); + } + + INFO("info: CU mask for the default stream is: 0x" << ss.str().c_str()); + if (globalCUMask.size() > 0) { + if (ss.str().compare(globalCUMask) != 0) { + INFO("Expected CU mask:" << globalCUMask.c_str() << + ", api returned:" << ss.str().c_str()); + REQUIRE(false); + } + } else { + for (int i = 0 ; i < min(cuMask.size(), defaultCUMask.size()); i++) { + if (cuMask[i] != defaultCUMask[i]) { + INFO("Expected CU mask " << defaultCUMask[i] << + ", api returned:" << cuMask[i]); + REQUIRE(false); + } + } + } + } + + SECTION("Verify with custom mask set") { + std::vector customMask(defaultCUMask); + hipStream_t stream; + customMask[0] = 0xe; + + HIP_CHECK(hipExtStreamCreateWithCUMask(&stream, customMask.size(), + customMask.data())); + ss.str(""); + for (int i = customMask.size() - 1; i >= 0; i--) { + ss << customMask[i]; + } + INFO("info: setting a custom CU mask 0x" << ss.str()); + + HIP_CHECK(hipExtStreamGetCUMask(stream, cuMask.size(), &cuMask[0])); + ss.str(""); + for (int i = cuMask.size() - 1; i >= 0; i--) { + ss << cuMask[i]; + } + + size_t found = ss.str().find_first_not_of("0"); + if (found != std::string::npos) { + ss.str(ss.str().substr(found, ss.str().length())); + } + + INFO("info: reading back CU mask 0x" << ss.str() << + " for stream " << stream); + + if (!gCUMask) { + for (size_t i = 0; i < customMask.size(); i++) { + if (customMask[i] != cuMask[i]) { + INFO("Error! expected CU mask:" << customMask[i] + << ", Received CU mask:" << cuMask[i]); + REQUIRE(false); + } + } + } + + HIP_CHECK(hipStreamDestroy(stream)); + } +} + +/** + * Negative tests for hipExtStreamGetCUMask. + */ +TEST_CASE("Unit_hipExtStreamGetCUMask_Negative") { + hipError_t ret; + constexpr int maxNum = 6; + std::vector cuMask(maxNum); + + SECTION("cuMask is nullptr") { + ret = hipExtStreamGetCUMask(0, cuMask.size(), nullptr); + REQUIRE(ret == hipErrorInvalidValue); + } + + SECTION("cuMaskSize is 0") { + ret = hipExtStreamGetCUMask(0, 0, &cuMask[0]); + REQUIRE(ret == hipErrorInvalidValue); + } +} diff --git a/projects/hip/tests/catch/unit/stream/hipStreamGetFlags.cc b/projects/hip/tests/catch/unit/stream/hipStreamGetFlags.cc index 583abe5068..9cbe37b6de 100644 --- a/projects/hip/tests/catch/unit/stream/hipStreamGetFlags.cc +++ b/projects/hip/tests/catch/unit/stream/hipStreamGetFlags.cc @@ -16,16 +16,24 @@ 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. */ +/** +Testcase Scenarios : +1) Test flag value of stream created with hipStreamCreateWithFlags/ + /hipStreamCreate/hipStreamCreateWithPriority. +2) Negative tests for hipStreamGetFlags api. +3) Test flag value when streams created with CUMask. +*/ + #include -TEST_CASE("Unit_hipStreamGetFlags_Negative") { - // Get flags for uninitialized stream + +/** + * Test flag value of stream created with various types. + */ +TEST_CASE("Unit_hipStreamGetFlags_BasicFunctionalities") { hipStream_t stream; - HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamDefault)); - REQUIRE(hipStreamGetFlags(stream, nullptr) == hipErrorInvalidValue); -} -TEST_CASE("Unit_hipStreamGetFlags") { - hipStream_t stream; - unsigned int flags; + unsigned int flags; + // Check flag value of stream created with hipStreamCreateWithFlags + SECTION("Check flag value of streams hipStreamCreateWithFlags") { HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamDefault)); HIP_CHECK(hipStreamGetFlags(stream, &flags)); REQUIRE(flags == hipStreamDefault); @@ -34,4 +42,50 @@ TEST_CASE("Unit_hipStreamGetFlags") { HIP_CHECK(hipStreamGetFlags(stream, &flags)); REQUIRE(flags == hipStreamNonBlocking); HIP_CHECK(hipStreamDestroy(stream)); + } + // Check flag value of stream created with hipStreamCreate + SECTION("Check flag value of streams hipStreamCreate") { + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamGetFlags(stream, &flags)); + REQUIRE(flags == hipStreamDefault); + HIP_CHECK(hipStreamDestroy(stream)); + } + // Check flag value of stream created with hipStreamCreateWithPriority + SECTION("Check flag value of streams hipStreamCreateWithPriority") { + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, 0)); + HIP_CHECK(hipStreamGetFlags(stream, &flags)); + REQUIRE(flags == hipStreamDefault); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, 0)); + HIP_CHECK(hipStreamGetFlags(stream, &flags)); + REQUIRE(flags == hipStreamNonBlocking); + HIP_CHECK(hipStreamDestroy(stream)); + } } + +/** + * Negative Scenarios + */ +TEST_CASE("Unit_hipStreamGetFlags_Negative") { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + // nullptr check + REQUIRE(hipStreamGetFlags(stream, nullptr) != hipSuccess); + REQUIRE(hipStreamGetFlags(nullptr, hipStreamDefault) != hipSuccess); + HIP_CHECK(hipStreamDestroy(stream)); +} + +#if HT_AMD +/** + * Test flag value when streams created with CUMask. + */ +TEST_CASE("Unit_hipStreamGetFlags_StreamsCreatedWithCUMask") { + hipStream_t stream; + unsigned int flags; + const uint32_t cuMask = 0xffffffff; + HIP_CHECK(hipExtStreamCreateWithCUMask(&stream, 1, &cuMask)); + HIP_CHECK(hipStreamGetFlags(stream, &flags)); + REQUIRE(flags == hipStreamDefault); + HIP_CHECK(hipStreamDestroy(stream)); +} +#endif diff --git a/projects/hip/tests/catch/unit/stream/hipStreamGetPriority.cc b/projects/hip/tests/catch/unit/stream/hipStreamGetPriority.cc index 3578fd9b5a..29f182ecdf 100644 --- a/projects/hip/tests/catch/unit/stream/hipStreamGetPriority.cc +++ b/projects/hip/tests/catch/unit/stream/hipStreamGetPriority.cc @@ -16,11 +16,28 @@ 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. */ +/** +Testcase Scenarios : +1) Negative tests for hipStreamGetPriority api. +2) Create stream and check default priority of stream is within range. +3) Create stream with high or low priority and check priority is set as expected. +4) Create stream with higher priority or lower priority for the priority range returned. +5) Create stream with CUMask and check priority is returned as expected. +*/ + #include + +/** + * Negative tests for hipStreamGetPriority api. + */ TEST_CASE("Unit_hipStreamGetPriority_Negative") { hipStream_t stream = 0; REQUIRE(hipStreamGetPriority(stream, nullptr) == hipErrorInvalidValue); } + +/** + * Create stream and check default priority of stream is within range. + */ TEST_CASE("Unit_hipStreamGetPriority_default") { int priority_low = 0; int priority_high = 0; @@ -37,6 +54,10 @@ TEST_CASE("Unit_hipStreamGetPriority_default") { REQUIRE(priority >= priority_high); HIP_CHECK(hipStreamDestroy(stream)); } + +/** + * Create stream with high priority and check priority is set as expected. + */ TEST_CASE("Unit_hipStreamGetPriority_high") { int priority_low = 0; int priority_high = 0; @@ -44,12 +65,17 @@ TEST_CASE("Unit_hipStreamGetPriority_high") { HIP_CHECK(hipSetDevice(devID)); HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); hipStream_t stream; - HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_high)); + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, + priority_high)); int priority = 0; HIP_CHECK(hipStreamGetPriority(stream, &priority)); REQUIRE(priority == priority_high); HIP_CHECK(hipStreamDestroy(stream)); } + +/** + * Create stream with higher priority for the priority range returned. + */ TEST_CASE("Unit_hipStreamGetPriority_higher") { int priority_low = 0; int priority_high = 0; @@ -57,12 +83,17 @@ TEST_CASE("Unit_hipStreamGetPriority_higher") { HIP_CHECK(hipSetDevice(devID)); HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); hipStream_t stream; - HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_high-1)); + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, + priority_high-1)); int priority = 0; HIP_CHECK(hipStreamGetPriority(stream, &priority)); REQUIRE(priority == priority_high); HIP_CHECK(hipStreamDestroy(stream)); } + +/** + * Create stream with low priority and check priority is set as expected. + */ TEST_CASE("Unit_hipStreamGetPriority_low") { int priority_low = 0; int priority_high = 0; @@ -70,12 +101,17 @@ TEST_CASE("Unit_hipStreamGetPriority_low") { HIP_CHECK(hipSetDevice(devID)); HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); hipStream_t stream; - HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_low)); + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, + priority_low)); int priority = 0; HIP_CHECK(hipStreamGetPriority(stream, &priority)); REQUIRE(priority_low == priority); HIP_CHECK(hipStreamDestroy(stream)); } + +/** + * Create stream with lower priority for the priority range returned. + */ TEST_CASE("Unit_hipStreamGetPriority_lower") { int priority_low = 0; int priority_high = 0; @@ -83,9 +119,35 @@ TEST_CASE("Unit_hipStreamGetPriority_lower") { HIP_CHECK(hipSetDevice(devID)); HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); hipStream_t stream; - HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_low+1)); + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, + priority_low+1)); int priority = 0; HIP_CHECK(hipStreamGetPriority(stream, &priority)); REQUIRE(priority_low == priority); HIP_CHECK(hipStreamDestroy(stream)); } + +#if HT_AMD +/** + * Create stream with CUMask and check priority is returned as expected. + */ +TEST_CASE("Unit_hipStreamGetPriority_StreamsWithCUMask") { + hipStream_t stream; + int priority = 0; + int priority_normal = 0; + int priority_low = 0; + int priority_high = 0; + // Test is to get the Stream Priority Range + HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); + priority_normal = priority_low + priority_high; + // Check if priorities are indeed supported + REQUIRE_FALSE(priority_low == priority_high); + // Creating a stream with hipExtStreamCreateWithCUMask and checking + // priority. + const uint32_t cuMask = 0xffffffff; + HIP_CHECK(hipExtStreamCreateWithCUMask(&stream, 1, &cuMask)); + HIP_CHECK(hipStreamGetPriority(stream, &priority)); + REQUIRE_FALSE(priority_normal != priority); + HIP_CHECK(hipStreamDestroy(stream)); +} +#endif