From 1bb4348c3b67c4b21982c4d963cc8659a99129bb Mon Sep 17 00:00:00 2001 From: milos-mozetic <118800401+milos-mozetic@users.noreply.github.com> Date: Sat, 8 Jul 2023 17:20:41 +0200 Subject: [PATCH] EXSWHTEC-224 - Test cases ID clean up and documentation for Execution Control (#85) - Test cases ID clean up and documentation for Execution Control - Update config_nvidia_linux_common.json - Disable test which failed on external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/85 [ROCm/hip-tests commit: e103027f81ba8923b4bf8d0c9f06e08d551953e2] --- .../config/config_amd_linux_common.json | 13 + .../config/config_amd_windows_common.json | 13 + .../config/config_nvidia_linux_common.json | 11 + .../config/config_nvidia_windows_common.json | 9 + .../catch/include/hip_test_defgroups.hh | 7 + projects/hip-tests/catch/unit/CMakeLists.txt | 1 + .../unit/executionControl/CMakeLists.txt | 11 + .../execution_control_common.cc | 27 ++ .../execution_control_common.hh | 25 ++ .../executionControl/hipFuncSetAttribute.cc | 260 ++++++++++++++++++ .../executionControl/hipFuncSetCacheConfig.cc | 105 +++++++ .../hipFuncSetSharedMemConfig.cc | 106 +++++++ 12 files changed, 588 insertions(+) create mode 100644 projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux_common.json create mode 100644 projects/hip-tests/catch/hipTestMain/config/config_nvidia_windows_common.json create mode 100644 projects/hip-tests/catch/unit/executionControl/CMakeLists.txt create mode 100644 projects/hip-tests/catch/unit/executionControl/execution_control_common.cc create mode 100644 projects/hip-tests/catch/unit/executionControl/execution_control_common.hh create mode 100644 projects/hip-tests/catch/unit/executionControl/hipFuncSetAttribute.cc create mode 100644 projects/hip-tests/catch/unit/executionControl/hipFuncSetCacheConfig.cc create mode 100644 projects/hip-tests/catch/unit/executionControl/hipFuncSetSharedMemConfig.cc diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json index b9c0e7c79f..9704e80d49 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json @@ -13,6 +13,19 @@ "Unit_hipInit_Negative", "Unit_hipDeviceReset_Positive_Basic", "Unit_hipDeviceReset_Positive_Threaded", + "Unit_hipFuncSetCacheConfig_Positive_Basic", + "Unit_hipFuncSetCacheConfig_Negative_Parameters", + "Unit_hipFuncSetSharedMemConfig_Positive_Basic", + "Unit_hipFuncSetSharedMemConfig_Negative_Parameters", + "Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize", + "Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout", + "Unit_hipFuncSetAttribute_Positive_Parameters", + "Unit_hipFuncSetAttribute_Negative_Parameters", + "NOTE: The following 4 tests are disabled due to defect - EXSWHTEC-240", + "Unit_hipFuncSetCacheConfig_Negative_Not_Supported", + "Unit_hipFuncSetSharedMemConfig_Negative_Not_Supported", + "Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize_Not_Supported", + "Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout_Not_Supported", "Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters", "Unit_hipGraphMemcpyNodeSetParamsToSymbol_Positive_Basic", "Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Positive_Basic", diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json index 6b7c976792..2d050a9ca4 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json @@ -95,6 +95,19 @@ "Unit_hipStreamSynchronize_NullStreamAndStreamPerThread", "Note: intermittent Seg fault failure ", "Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags", + "Unit_hipFuncSetCacheConfig_Positive_Basic", + "Unit_hipFuncSetCacheConfig_Negative_Parameters", + "Unit_hipFuncSetSharedMemConfig_Positive_Basic", + "Unit_hipFuncSetSharedMemConfig_Negative_Parameters", + "Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize", + "Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout", + "Unit_hipFuncSetAttribute_Positive_Parameters", + "Unit_hipFuncSetAttribute_Negative_Parameters", + "NOTE: The following 4 tests are disabled due to defect - EXSWHTEC-240", + "Unit_hipFuncSetCacheConfig_Negative_Not_Supported", + "Unit_hipFuncSetSharedMemConfig_Negative_Not_Supported", + "Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize_Not_Supported", + "Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout_Not_Supported", "Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters", "Unit_hipGraphMemcpyNodeSetParamsToSymbol_Positive_Basic", "Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Positive_Basic", diff --git a/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux_common.json b/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux_common.json new file mode 100644 index 0000000000..7766cbe28d --- /dev/null +++ b/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux_common.json @@ -0,0 +1,11 @@ +{ + "DisabledTests": + [ + "Unit_hipFuncSetCacheConfig_Negative_Not_Supported", + "Unit_hipFuncSetSharedMemConfig_Negative_Not_Supported", + "Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize_Not_Supported", + "Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout_Not_Supported", + "=== Below test fails in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/85 ===", + "Unit_hipFuncSetAttribute_Negative_Parameters" + ] +} diff --git a/projects/hip-tests/catch/hipTestMain/config/config_nvidia_windows_common.json b/projects/hip-tests/catch/hipTestMain/config/config_nvidia_windows_common.json new file mode 100644 index 0000000000..8247b49e3a --- /dev/null +++ b/projects/hip-tests/catch/hipTestMain/config/config_nvidia_windows_common.json @@ -0,0 +1,9 @@ +{ + "DisabledTests": + [ + "Unit_hipFuncSetCacheConfig_Negative_Not_Supported", + "Unit_hipFuncSetSharedMemConfig_Negative_Not_Supported", + "Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize_Not_Supported", + "Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout_Not_Supported" + ] +} diff --git a/projects/hip-tests/catch/include/hip_test_defgroups.hh b/projects/hip-tests/catch/include/hip_test_defgroups.hh index ce6158a72d..172f228f84 100644 --- a/projects/hip-tests/catch/include/hip_test_defgroups.hh +++ b/projects/hip-tests/catch/include/hip_test_defgroups.hh @@ -72,6 +72,13 @@ THE SOFTWARE. * @} */ +/** + * @defgroup ExecutionTest Execution Control + * @{ + * This section describes tests for the execution control functions of HIP runtime API. + * @} + */ + /** * @defgroup ShflTest warp shuffle function Management * @{ diff --git a/projects/hip-tests/catch/unit/CMakeLists.txt b/projects/hip-tests/catch/unit/CMakeLists.txt index c26af21b02..06739c623a 100644 --- a/projects/hip-tests/catch/unit/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/CMakeLists.txt @@ -42,3 +42,4 @@ add_subdirectory(callback) # Vulkan interop APIs currently undefined for Nvidia add_subdirectory(vulkan_interop) endif() +add_subdirectory(executionControl) diff --git a/projects/hip-tests/catch/unit/executionControl/CMakeLists.txt b/projects/hip-tests/catch/unit/executionControl/CMakeLists.txt new file mode 100644 index 0000000000..4c28438009 --- /dev/null +++ b/projects/hip-tests/catch/unit/executionControl/CMakeLists.txt @@ -0,0 +1,11 @@ +set(TEST_SRC + execution_control_common.cc + hipFuncSetCacheConfig.cc + hipFuncSetSharedMemConfig.cc + hipFuncSetAttribute.cc +) + +hip_add_exe_to_target(NAME ExecutionControlTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests + COMPILE_OPTIONS -std=c++17) \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/executionControl/execution_control_common.cc b/projects/hip-tests/catch/unit/executionControl/execution_control_common.cc new file mode 100644 index 0000000000..f9a2cc1675 --- /dev/null +++ b/projects/hip-tests/catch/unit/executionControl/execution_control_common.cc @@ -0,0 +1,27 @@ +/* +Copyright (c) 2022 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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS 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 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. +*/ + +#include "execution_control_common.hh" + +#include + +__global__ void kernel() {} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/executionControl/execution_control_common.hh b/projects/hip-tests/catch/unit/executionControl/execution_control_common.hh new file mode 100644 index 0000000000..e66fcf28f3 --- /dev/null +++ b/projects/hip-tests/catch/unit/executionControl/execution_control_common.hh @@ -0,0 +1,25 @@ +/* +Copyright (c) 2022 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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS 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 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. +*/ + +#pragma once + +__global__ void kernel(); \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/executionControl/hipFuncSetAttribute.cc b/projects/hip-tests/catch/unit/executionControl/hipFuncSetAttribute.cc new file mode 100644 index 0000000000..785cc91e45 --- /dev/null +++ b/projects/hip-tests/catch/unit/executionControl/hipFuncSetAttribute.cc @@ -0,0 +1,260 @@ +/* +Copyright (c) 2022 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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS 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 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. +*/ + +#include "execution_control_common.hh" + +#include +#include + +/** + * @addtogroup hipFuncSetAttribute hipFuncSetAttribute + * @{ + * @ingroup ExecutionTest + * `hipFuncSetAttribute(const void* func, hipFuncAttribute attr, int value)` - + * Set attribute for a specific function. + */ + +/** + * Test Description + * ------------------------ + * - Sets maximum dynamic shared memory size to the non-default value. + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/executionControl/hipFuncSetAttribute.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize") { + HIP_CHECK(hipFuncSetAttribute(reinterpret_cast(kernel), + hipFuncAttributeMaxDynamicSharedMemorySize, 1024)); + + hipFuncAttributes attributes; + HIP_CHECK(hipFuncGetAttributes(&attributes, reinterpret_cast(kernel))); + + REQUIRE(attributes.maxDynamicSharedSizeBytes == 1024); +} + +/** + * Test Description + * ------------------------ + * - Sets preferred shared memory carveout to the non-default value. + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/executionControl/hipFuncSetAttribute.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout") { + HIP_CHECK(hipFuncSetAttribute(reinterpret_cast(kernel), + hipFuncAttributePreferredSharedMemoryCarveout, 50)); + + hipFuncAttributes attributes; + HIP_CHECK(hipFuncGetAttributes(&attributes, reinterpret_cast(kernel))); + + REQUIRE(attributes.preferredShmemCarveout == 50); +} + +/** + * Test Description + * ------------------------ + * - Validates handling of valid arguments: + * -# When `hipFuncAttributeMaxDynamicSharedMemorySize == 0` + * - Expected output: return `hipSuccess` + * -# When `hipFuncAttributeMaxDynamicSharedMemorySize == maxSharedMemoryPerBlock - sharedSizeBytes` + * - Expected output: return `hipSuccess` + * -# When `hipFuncAttributePreferredSharedMemoryCarveout` is 0% + * - Expected output: return `hipSuccess` + * -# When `hipFuncAttributePreferredSharedMemoryCarveout` is 100% + * - Expected output: return `hipSuccess` + * -# When `hipFuncAttributePreferredSharedMemoryCarveout` is default (-1) + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/executionControl/hipFuncSetAttribute.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipFuncSetAttribute_Positive_Parameters") { + SECTION("hipFuncAttributeMaxDynamicSharedMemorySize == 0") { + HIP_CHECK(hipFuncSetAttribute(reinterpret_cast(kernel), + hipFuncAttributeMaxDynamicSharedMemorySize, 0)); + } + + SECTION( + "hipFuncAttributeMaxDynamicSharedMemorySize == maxSharedMemoryPerBlock - sharedSizeBytes") { + // The sum of this value and the function attribute sharedSizeBytes cannot exceed the device + // attribute cudaDevAttrMaxSharedMemoryPerBlockOptin + int max_shared; + HIP_CHECK(hipDeviceGetAttribute(&max_shared, hipDeviceAttributeMaxSharedMemoryPerBlock, 0)); + + hipFuncAttributes attributes; + HIP_CHECK(hipFuncGetAttributes(&attributes, reinterpret_cast(kernel))); + + HIP_CHECK(hipFuncSetAttribute(reinterpret_cast(kernel), + hipFuncAttributeMaxDynamicSharedMemorySize, + max_shared - attributes.sharedSizeBytes)); + } + + SECTION("hipFuncAttributePreferredSharedMemoryCarveout == 0") { + HIP_CHECK(hipFuncSetAttribute(reinterpret_cast(kernel), + hipFuncAttributePreferredSharedMemoryCarveout, 0)); + } + + SECTION("hipFuncAttributePreferredSharedMemoryCarveout == 100") { + HIP_CHECK(hipFuncSetAttribute(reinterpret_cast(kernel), + hipFuncAttributePreferredSharedMemoryCarveout, 100)); + } + + SECTION("hipFuncAttributePreferredSharedMemoryCarveout == -1 (default)") { + HIP_CHECK(hipFuncSetAttribute(reinterpret_cast(kernel), + hipFuncAttributePreferredSharedMemoryCarveout, -1)); + } +} + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When pointer to the kernel function is `nullptr` + * - Expected output: return `hipErrorInvalidDeviceFunction` + * -# When the attribute is invalid + * - Expected output: return `hipErrorInvalidValue` + * -# When `hipFuncAttributeMaxDynamicSharedMemorySize < 0` + * - Expected output: return `hipErrorInvalidValue` + * -# When `hipFuncAttributeMaxDynamicSharedMemorySize > maxSharedMemoryPerBlock - sharedSizeBytes` + * - Expected output: return `hipErrorInvalidValue` + * -# When `hipFuncAttributePreferredSharedMemoryCarveout` is negative + * - Expected output: return `hipErrorInvalidValue` + * -# When `hipFuncAttributePreferredSharedMemoryCarveout` is above 100% + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/executionControl/hipFuncSetAttribute.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipFuncSetAttribute_Negative_Parameters") { + SECTION("func == nullptr") { + HIP_CHECK_ERROR(hipFuncSetAttribute(nullptr, hipFuncAttributePreferredSharedMemoryCarveout, 50), + hipErrorInvalidDeviceFunction); + } + + SECTION("invalid attribute") { + HIP_CHECK_ERROR( + hipFuncSetAttribute(reinterpret_cast(kernel), static_cast(-1), 50), + hipErrorInvalidValue); + } + + SECTION("hipFuncAttributeMaxDynamicSharedMemorySize < 0") { + HIP_CHECK_ERROR(hipFuncSetAttribute(reinterpret_cast(kernel), + hipFuncAttributeMaxDynamicSharedMemorySize, -1), + hipErrorInvalidValue); + } + + SECTION( + "hipFuncAttributeMaxDynamicSharedMemorySize > maxSharedMemoryPerBlock - sharedSizeBytes") { + // The sum of this value and the function attribute sharedSizeBytes cannot exceed the device + // attribute cudaDevAttrMaxSharedMemoryPerBlockOptin + int max_shared; + HIP_CHECK(hipDeviceGetAttribute(&max_shared, hipDeviceAttributeMaxSharedMemoryPerBlock, 0)); + + hipFuncAttributes attributes; + HIP_CHECK(hipFuncGetAttributes(&attributes, reinterpret_cast(kernel))); + + HIP_CHECK_ERROR(hipFuncSetAttribute(reinterpret_cast(kernel), + hipFuncAttributeMaxDynamicSharedMemorySize, + max_shared - attributes.sharedSizeBytes + 1), + hipErrorInvalidValue); + } + + SECTION("hipFuncAttributePreferredSharedMemoryCarveout < -1") { + HIP_CHECK_ERROR(hipFuncSetAttribute(reinterpret_cast(kernel), + hipFuncAttributePreferredSharedMemoryCarveout, -2), + hipErrorInvalidValue); + } + + SECTION("hipFuncAttributePreferredSharedMemoryCarveout > 100") { + HIP_CHECK_ERROR(hipFuncSetAttribute(reinterpret_cast(kernel), + hipFuncAttributePreferredSharedMemoryCarveout, 101), + hipErrorInvalidValue); + } +} + +/** + * Test Description + * ------------------------ + * - Sets `hipFuncAttributeMaxDynamicSharedMemorySize` to the non-supported value + * - Expected output: return `hipErrorNotSupported` + * Test source + * ------------------------ + * - unit/executionControl/hipFuncSetAttribute.cc + * Test requirements + * ------------------------ + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize_Not_Supported") { + hipFuncAttributes old_attributes; + HIP_CHECK(hipFuncGetAttributes(&old_attributes, reinterpret_cast(kernel))); + + HIP_CHECK_ERROR(hipFuncSetAttribute(reinterpret_cast(kernel), + hipFuncAttributeMaxDynamicSharedMemorySize, 1024), + hipErrorNotSupported); + + hipFuncAttributes new_attributes; + HIP_CHECK(hipFuncGetAttributes(&new_attributes, reinterpret_cast(kernel))); + + REQUIRE(old_attributes.maxDynamicSharedSizeBytes == new_attributes.maxDynamicSharedSizeBytes); +} + +/** + * Test Description + * ------------------------ + * - Sets `hipFuncAttributePreferredSharedMemoryCarveout` to the non-supported value + * - Expected output: return `hipErrorNotSupported` + * Test source + * ------------------------ + * - unit/executionControl/hipFuncSetAttribute.cc + * Test requirements + * ------------------------ + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout_Not_Supported") { + hipFuncAttributes old_attributes; + HIP_CHECK(hipFuncGetAttributes(&old_attributes, reinterpret_cast(kernel))); + + HIP_CHECK_ERROR(hipFuncSetAttribute(reinterpret_cast(kernel), + hipFuncAttributePreferredSharedMemoryCarveout, 50), + hipErrorNotSupported); + + hipFuncAttributes new_attributes; + HIP_CHECK(hipFuncGetAttributes(&new_attributes, reinterpret_cast(kernel))); + + REQUIRE(old_attributes.preferredShmemCarveout == new_attributes.preferredShmemCarveout); +} diff --git a/projects/hip-tests/catch/unit/executionControl/hipFuncSetCacheConfig.cc b/projects/hip-tests/catch/unit/executionControl/hipFuncSetCacheConfig.cc new file mode 100644 index 0000000000..1f5ce8dfda --- /dev/null +++ b/projects/hip-tests/catch/unit/executionControl/hipFuncSetCacheConfig.cc @@ -0,0 +1,105 @@ +/* +Copyright (c) 2022 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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS 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 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. +*/ + +#include "execution_control_common.hh" + +#include +#include + +/** + * @addtogroup hipFuncSetCacheConfig hipFuncSetCacheConfig + * @{ + * @ingroup ExecutionTest + * `hipFuncSetCacheConfig(const void* func, hipFuncCache_t config)` - + * Set Cache configuration for a specific function. + */ + +namespace { +constexpr std::array kCacheConfigs{ + hipFuncCachePreferNone, hipFuncCachePreferShared, hipFuncCachePreferL1, + hipFuncCachePreferEqual}; +} // anonymous namespace + +/** + * Test Description + * ------------------------ + * - Sets cache config for every cache config enumeration. + * Test source + * ------------------------ + * - unit/executionControl/hipFuncSetCacheConfig.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipFuncSetCacheConfig_Positive_Basic") { + const auto cache_config = GENERATE(from_range(begin(kCacheConfigs), end(kCacheConfigs))); + + HIP_CHECK(hipFuncSetCacheConfig(reinterpret_cast(kernel), cache_config)); + + kernel<<<1, 1>>>(); + HIP_CHECK(hipDeviceSynchronize()); +} + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When pointer to the kernel function is `nullptr` + * - Expected output: return `hipErrorInvalidDeviceFunction` + * -# When cache config enumeration is invalid + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/executionControl/hipFuncSetCacheConfig.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipFuncSetCacheConfig_Negative_Parameters") { + SECTION("func == nullptr") { + HIP_CHECK_ERROR(hipFuncSetCacheConfig(nullptr, hipFuncCachePreferNone), + hipErrorInvalidDeviceFunction); + } + SECTION("invalid cache config") { + HIP_CHECK_ERROR( + hipFuncSetCacheConfig(reinterpret_cast(kernel), static_cast(-1)), + hipErrorInvalidValue); + } +} + +/** + * Test Description + * ------------------------ + * - Sets cache config that is not supported. + * - Expected output: return `hipErrorNotSupported` + * Test source + * ------------------------ + * - unit/executionControl/hipFuncSetCacheConfig.cc + * Test requirements + * ------------------------ + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipFuncSetCacheConfig_Negative_Not_Supported") { + HIP_CHECK_ERROR(hipFuncSetCacheConfig(reinterpret_cast(kernel), hipFuncCachePreferNone), + hipErrorNotSupported); +} diff --git a/projects/hip-tests/catch/unit/executionControl/hipFuncSetSharedMemConfig.cc b/projects/hip-tests/catch/unit/executionControl/hipFuncSetSharedMemConfig.cc new file mode 100644 index 0000000000..c44a83b3be --- /dev/null +++ b/projects/hip-tests/catch/unit/executionControl/hipFuncSetSharedMemConfig.cc @@ -0,0 +1,106 @@ +/* +Copyright (c) 2022 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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS 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 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. +*/ + +#include "execution_control_common.hh" + +#include +#include + +/** + * @addtogroup hipFuncSetSharedMemConfig hipFuncSetSharedMemConfig + * @{ + * @ingroup ExecutionTest + * `hipFuncSetSharedMemConfig(const void* func, hipSharedMemConfig config)` - + * Set shared memory configuation for a specific function. + */ + +namespace { +constexpr std::array kSharedMemConfigs{ + hipSharedMemBankSizeDefault, hipSharedMemBankSizeFourByte, hipSharedMemBankSizeEightByte}; +} // anonymous namespace + +/** + * Test Description + * ------------------------ + * - Sets cache config for every shared memory config enumeration. + * Test source + * ------------------------ + * - unit/executionControl/hipFuncSetSharedMemConfig.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipFuncSetSharedMemConfig_Positive_Basic") { + const auto shared_mem_config = + GENERATE(from_range(begin(kSharedMemConfigs), end(kSharedMemConfigs))); + + HIP_CHECK(hipFuncSetSharedMemConfig(reinterpret_cast(kernel), shared_mem_config)); + + kernel<<<1, 1>>>(); + HIP_CHECK(hipDeviceSynchronize()); +} + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When pointer to the kernel function is `nullptr` + * - Expected output: return `hipErrorInvalidDeviceFunction` + * -# When shared memory config enumeration is invalid + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/executionControl/hipFuncSetSharedMemConfig.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipFuncSetSharedMemConfig_Negative_Parameters") { + SECTION("func == nullptr") { + HIP_CHECK_ERROR(hipFuncSetSharedMemConfig(nullptr, hipSharedMemBankSizeDefault), + hipErrorInvalidDeviceFunction); + } + SECTION("invalid shared mem config") { + HIP_CHECK_ERROR(hipFuncSetSharedMemConfig(reinterpret_cast(kernel), + static_cast(-1)), + hipErrorInvalidValue); + } +} + +/** + * Test Description + * ------------------------ + * - Sets shared memory config that is not supported. + * - Expected output: return `hipErrorNotSupported` + * Test source + * ------------------------ + * - unit/executionControl/hipFuncSetSharedMemConfig.cc + * Test requirements + * ------------------------ + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipFuncSetSharedMemConfig_Negative_Not_Supported") { + HIP_CHECK_ERROR( + hipFuncSetSharedMemConfig(reinterpret_cast(kernel), hipSharedMemBankSizeDefault), + hipErrorNotSupported); +}