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: e103027f81]
Этот коммит содержится в:
milos-mozetic
коммит произвёл
GitHub
GitHub
родитель
adaf3dc14e
Коммит
1bb4348c3b
@@ -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",
|
||||
|
||||
@@ -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",
|
||||
|
||||
@@ -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"
|
||||
]
|
||||
}
|
||||
@@ -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"
|
||||
]
|
||||
}
|
||||
@@ -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
|
||||
* @{
|
||||
|
||||
@@ -42,3 +42,4 @@ add_subdirectory(callback)
|
||||
# Vulkan interop APIs currently undefined for Nvidia
|
||||
add_subdirectory(vulkan_interop)
|
||||
endif()
|
||||
add_subdirectory(executionControl)
|
||||
|
||||
@@ -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)
|
||||
@@ -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 <hip_test_common.hh>
|
||||
|
||||
__global__ void kernel() {}
|
||||
@@ -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();
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_runtime_api.h>
|
||||
|
||||
/**
|
||||
* @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<void*>(kernel),
|
||||
hipFuncAttributeMaxDynamicSharedMemorySize, 1024));
|
||||
|
||||
hipFuncAttributes attributes;
|
||||
HIP_CHECK(hipFuncGetAttributes(&attributes, reinterpret_cast<void*>(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<void*>(kernel),
|
||||
hipFuncAttributePreferredSharedMemoryCarveout, 50));
|
||||
|
||||
hipFuncAttributes attributes;
|
||||
HIP_CHECK(hipFuncGetAttributes(&attributes, reinterpret_cast<void*>(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<void*>(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<void*>(kernel)));
|
||||
|
||||
HIP_CHECK(hipFuncSetAttribute(reinterpret_cast<void*>(kernel),
|
||||
hipFuncAttributeMaxDynamicSharedMemorySize,
|
||||
max_shared - attributes.sharedSizeBytes));
|
||||
}
|
||||
|
||||
SECTION("hipFuncAttributePreferredSharedMemoryCarveout == 0") {
|
||||
HIP_CHECK(hipFuncSetAttribute(reinterpret_cast<void*>(kernel),
|
||||
hipFuncAttributePreferredSharedMemoryCarveout, 0));
|
||||
}
|
||||
|
||||
SECTION("hipFuncAttributePreferredSharedMemoryCarveout == 100") {
|
||||
HIP_CHECK(hipFuncSetAttribute(reinterpret_cast<void*>(kernel),
|
||||
hipFuncAttributePreferredSharedMemoryCarveout, 100));
|
||||
}
|
||||
|
||||
SECTION("hipFuncAttributePreferredSharedMemoryCarveout == -1 (default)") {
|
||||
HIP_CHECK(hipFuncSetAttribute(reinterpret_cast<void*>(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<void*>(kernel), static_cast<hipFuncAttribute>(-1), 50),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("hipFuncAttributeMaxDynamicSharedMemorySize < 0") {
|
||||
HIP_CHECK_ERROR(hipFuncSetAttribute(reinterpret_cast<void*>(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<void*>(kernel)));
|
||||
|
||||
HIP_CHECK_ERROR(hipFuncSetAttribute(reinterpret_cast<void*>(kernel),
|
||||
hipFuncAttributeMaxDynamicSharedMemorySize,
|
||||
max_shared - attributes.sharedSizeBytes + 1),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("hipFuncAttributePreferredSharedMemoryCarveout < -1") {
|
||||
HIP_CHECK_ERROR(hipFuncSetAttribute(reinterpret_cast<void*>(kernel),
|
||||
hipFuncAttributePreferredSharedMemoryCarveout, -2),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("hipFuncAttributePreferredSharedMemoryCarveout > 100") {
|
||||
HIP_CHECK_ERROR(hipFuncSetAttribute(reinterpret_cast<void*>(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<void*>(kernel)));
|
||||
|
||||
HIP_CHECK_ERROR(hipFuncSetAttribute(reinterpret_cast<void*>(kernel),
|
||||
hipFuncAttributeMaxDynamicSharedMemorySize, 1024),
|
||||
hipErrorNotSupported);
|
||||
|
||||
hipFuncAttributes new_attributes;
|
||||
HIP_CHECK(hipFuncGetAttributes(&new_attributes, reinterpret_cast<void*>(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<void*>(kernel)));
|
||||
|
||||
HIP_CHECK_ERROR(hipFuncSetAttribute(reinterpret_cast<void*>(kernel),
|
||||
hipFuncAttributePreferredSharedMemoryCarveout, 50),
|
||||
hipErrorNotSupported);
|
||||
|
||||
hipFuncAttributes new_attributes;
|
||||
HIP_CHECK(hipFuncGetAttributes(&new_attributes, reinterpret_cast<void*>(kernel)));
|
||||
|
||||
REQUIRE(old_attributes.preferredShmemCarveout == new_attributes.preferredShmemCarveout);
|
||||
}
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_runtime_api.h>
|
||||
|
||||
/**
|
||||
* @addtogroup hipFuncSetCacheConfig hipFuncSetCacheConfig
|
||||
* @{
|
||||
* @ingroup ExecutionTest
|
||||
* `hipFuncSetCacheConfig(const void* func, hipFuncCache_t config)` -
|
||||
* Set Cache configuration for a specific function.
|
||||
*/
|
||||
|
||||
namespace {
|
||||
constexpr std::array<hipFuncCache_t, 4> 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<void*>(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<void*>(kernel), static_cast<hipFuncCache_t>(-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<void*>(kernel), hipFuncCachePreferNone),
|
||||
hipErrorNotSupported);
|
||||
}
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_runtime_api.h>
|
||||
|
||||
/**
|
||||
* @addtogroup hipFuncSetSharedMemConfig hipFuncSetSharedMemConfig
|
||||
* @{
|
||||
* @ingroup ExecutionTest
|
||||
* `hipFuncSetSharedMemConfig(const void* func, hipSharedMemConfig config)` -
|
||||
* Set shared memory configuation for a specific function.
|
||||
*/
|
||||
|
||||
namespace {
|
||||
constexpr std::array<hipSharedMemConfig, 3> 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<void*>(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<void*>(kernel),
|
||||
static_cast<hipSharedMemConfig>(-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<void*>(kernel), hipSharedMemBankSizeDefault),
|
||||
hipErrorNotSupported);
|
||||
}
|
||||
Ссылка в новой задаче
Block a user