diff --git a/projects/hip-tests/catch/unit/stream/CMakeLists.txt b/projects/hip-tests/catch/unit/stream/CMakeLists.txt index a1980e4d18..bdaa5050c9 100644 --- a/projects/hip-tests/catch/unit/stream/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/stream/CMakeLists.txt @@ -20,7 +20,8 @@ set(TEST_SRC hipStreamGetDevice.cc hipStreamLegacy_Ext.cc hipStreamLegacy_compiler_options.cc - hipStreamGetId.cc) + hipStreamGetId.cc + hipStreamSetGetAttributes.cc) if(HIP_PLATFORM MATCHES "amd") set(TEST_SRC ${TEST_SRC} hipStreamGetCUMask.cc hipStreamWithCUMask.cc diff --git a/projects/hip-tests/catch/unit/stream/hipStreamSetGetAttributes.cc b/projects/hip-tests/catch/unit/stream/hipStreamSetGetAttributes.cc new file mode 100644 index 0000000000..bd42019e77 --- /dev/null +++ b/projects/hip-tests/catch/unit/stream/hipStreamSetGetAttributes.cc @@ -0,0 +1,183 @@ +/* +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING 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 ANY 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 + +/** + * Kernel to double each element in the given array + */ +static __global__ void doubleKernel(int *arr) { + size_t i = threadIdx.x; + arr[i] += arr[i]; +} + +/** + * Test Description + * ------------------------ + * - This test case checks the following scenario, + * - 1) Using hipStreamGetAttribute, get the default value of attribute + * - hipStreamAttributeSynchronizationPolicy + * - 2) Using hipStreamSetAttribute, set the possible values of + hipStreamAttributeSynchronizationPolicy attribute + * - 3) And validate it by getting value using hipStreamGetAttribute + * Test source + * ------------------------ + * - unit/stream/hipStreamSetGetAttributes.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipStreamSetAttribute_hipStreamGetAttribute_Basic") { + constexpr int N = 1024; + int hostMem[N]; + for (int i = 0; i < N; i++) { + hostMem[i] = 5; + } + + int *devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, N * sizeof(int))); + REQUIRE(devMem != nullptr); + + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + + hipStreamAttrID attr = hipStreamAttributeSynchronizationPolicy; + hipStreamAttrValue valueToSet; + hipSynchronizationPolicy syncPolicy = + GENERATE(hipSynchronizationPolicy::hipSyncPolicyAuto, + hipSynchronizationPolicy::hipSyncPolicySpin, + hipSynchronizationPolicy::hipSyncPolicyYield, + hipSynchronizationPolicy::hipSyncPolicyBlockingSync); + valueToSet.syncPolicy = syncPolicy; + HIP_CHECK(hipStreamSetAttribute(stream, attr, &valueToSet)); + + hipStreamAttrValue valueOut; + HIP_CHECK(hipStreamGetAttribute(stream, attr, &valueOut)); + REQUIRE(valueOut.syncPolicy == syncPolicy); + + HIP_CHECK(hipMemcpyAsync(devMem, hostMem, N * sizeof(int), + hipMemcpyHostToDevice, stream)); + doubleKernel<<<1, N, 0, stream>>>(devMem); + HIP_CHECK(hipMemcpyAsync(hostMem, devMem, N * sizeof(int), + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + for (int i = 0; i < N; i++) { + INFO("At index " << i << " Expected value = 10 " + << " Got value = " << hostMem[i]); + REQUIRE(hostMem[i] == 10); + } + + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(devMem)); +} + +/** + * Test Description + * ------------------------ + * - This test case checks the following Negative scenarios + * - for the hipStreamGetAttribute, + * - 1) With Invalid stream + * - 2) With Invalid attribute + * - 3) With Invalid Attribute value + * Test source + * ------------------------ + * - unit/stream/hipStreamSetGetAttributes.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipStreamGetAttribute_Negative") { + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + + hipStreamAttrID attr = hipStreamAttributeSynchronizationPolicy; + hipStreamAttrValue valueOut; + + SECTION("With Invalid stream") { + hipStream_t invalidStream = reinterpret_cast(-1); + HIP_CHECK_ERROR(hipStreamGetAttribute(invalidStream, attr, &valueOut), + hipErrorInvalidResourceHandle); + } + + SECTION("With invalid attribute") { + attr = static_cast(-1); + HIP_CHECK_ERROR(hipStreamGetAttribute(stream, attr, &valueOut), + hipErrorInvalidValue); + } +#if HT_AMD + // Facing segmentation fault issue in CUDA + SECTION("With invalid attribute value") { + HIP_CHECK_ERROR(hipStreamGetAttribute(stream, attr, nullptr), + hipErrorInvalidValue); + } +#endif + HIP_CHECK(hipStreamDestroy(stream)); +} + +/** + * Test Description + * ------------------------ + * - This test case checks the following Negative scenarios + * - for the hipStreamSetAttribute, + * - 1) With Invalid stream + * - 2) With Invalid attribute + * - 3) With Invalid Attribute value + * - 4) With Invalid hipSynchronizationPolicy + * Test source + * ------------------------ + * - unit/stream/hipStreamSetGetAttributes.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipStreamSetAttribute_Negative") { + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + + hipStreamAttrID attr = hipStreamAttributeSynchronizationPolicy; + hipStreamAttrValue valueToSet; + + SECTION("With Invalid stream") { + hipStream_t invalidStream = reinterpret_cast(-1); + valueToSet.syncPolicy = hipSynchronizationPolicy::hipSyncPolicyAuto; + HIP_CHECK_ERROR(hipStreamSetAttribute(invalidStream, attr, &valueToSet), + hipErrorInvalidResourceHandle); + } + + SECTION("With invalid attribute") { + attr = static_cast(-1); + HIP_CHECK_ERROR(hipStreamSetAttribute(stream, attr, &valueToSet), + hipErrorInvalidValue); + } +#if HT_AMD + // Facing segmentation fault issue in CUDA + SECTION("With invalid attribute value") { + HIP_CHECK_ERROR(hipStreamSetAttribute(stream, attr, nullptr), + hipErrorInvalidValue); + } +#endif + SECTION("With Invalid hipSynchronizationPolicy") { + valueToSet.syncPolicy = static_cast(-1); + HIP_CHECK_ERROR(hipStreamSetAttribute(stream, attr, &valueToSet), + hipErrorInvalidValue); + } + + HIP_CHECK(hipStreamDestroy(stream)); +}