SWDEV-1 - Merge github PRs to amd-staging
- https://github.com/ROCm/hip-tests/pull/119
- https://github.com/ROCm/hip-tests/pull/151
- https://github.com/ROCm/hip-tests/pull/57
- https://github.com/ROCm/hip-tests/pull/58
- https://github.com/ROCm/hip-tests/pull/59
- https://github.com/ROCm/hip-tests/pull/60
- https://github.com/ROCm/hip-tests/pull/99
- https://github.com/ROCm/hip-tests/pull/139
- https://github.com/ROCm/hip-tests/pull/152
- https://github.com/ROCm/hip-tests/pull/48
- https://github.com/ROCm/hip-tests/pull/54
- https://github.com/ROCm/hip-tests/pull/53
- https://github.com/ROCm/hip-tests/pull/24
- https://github.com/ROCm/hip-tests/pull/23
- https://github.com/ROCm/hip-tests/pull/22
- https://github.com/ROCm/hip-tests/pull/21
- https://github.com/ROCm/hip-tests/pull/20
- https://github.com/ROCm/hip-tests/pull/14
- https://github.com/ROCm/hip-tests/pull/8
Change-Id: I1eea54cd1436f3ddbfd5c1b3b2f672eb81d03cd4
[ROCm/hip-tests commit: 96df1fde80]
This commit is contained in:
Rakesh Roy
@@ -18,6 +18,9 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
add_subdirectory(memset)
|
||||
add_subdirectory(memcpy)
|
||||
add_subdirectory(kernelLaunch)
|
||||
add_subdirectory(stream)
|
||||
add_subdirectory(event)
|
||||
add_subdirectory(example)
|
||||
|
||||
@@ -0,0 +1,37 @@
|
||||
# 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.
|
||||
|
||||
set(TEST_SRC
|
||||
kernel_launch_common.cc
|
||||
triple_chevron.cc
|
||||
hipLaunchKernel.cc
|
||||
hipLaunchCooperativeKernel.cc
|
||||
)
|
||||
|
||||
if(HIP_PLATFORM MATCHES "amd")
|
||||
set(TEST_SRC ${TEST_SRC}
|
||||
hipExtLaunchKernel.cc
|
||||
)
|
||||
endif()
|
||||
|
||||
hip_add_exe_to_target(NAME KernelLaunchPerformance
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME build_tests
|
||||
COMPILE_OPTIONS -std=c++17)
|
||||
@@ -0,0 +1,120 @@
|
||||
/*
|
||||
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 "kernel_launch_common.hh"
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
/**
|
||||
* @addtogroup kernelLaunch kernel launch
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
* Contains performance tests for kernel launch overhead benchmarking.
|
||||
*/
|
||||
|
||||
template <KernelType kernel_type, bool timer_type>
|
||||
class ExtLaunchKernelBenchmark
|
||||
: public KernelLaunchBenchmark<ExtLaunchKernelBenchmark<kernel_type, timer_type>, timer_type> {
|
||||
public:
|
||||
constexpr void LaunchKernel() {
|
||||
if constexpr (kernel_type == KernelType::kNull) {
|
||||
error_ = hipExtLaunchKernel(reinterpret_cast<void*>(NullKernel), 1, 1, nullptr, 0, nullptr,
|
||||
events_[0], events_[1], 0u);
|
||||
} else if constexpr (kernel_type == KernelType::kSmall) {
|
||||
error_ = hipExtLaunchKernel(reinterpret_cast<void*>(KernelWithSmallArgs), 1, 1,
|
||||
small_kernel_args_, 0, nullptr, events_[0], events_[1], 0u);
|
||||
} else if constexpr (kernel_type == KernelType::kMedium) {
|
||||
error_ = hipExtLaunchKernel(reinterpret_cast<void*>(KernelWithMediumArgs), 1, 1,
|
||||
medium_kernel_args_, 0, nullptr, events_[0], events_[1], 0u);
|
||||
} else if constexpr (kernel_type == KernelType::kLarge) {
|
||||
error_ = hipExtLaunchKernel(reinterpret_cast<void*>(KernelWithLargeArgs), 1, 1,
|
||||
large_kernel_args_, 0, nullptr, events_[0], events_[1], 0u);
|
||||
} else
|
||||
;
|
||||
}
|
||||
|
||||
hipError_t GetError() { return error_; }
|
||||
|
||||
private:
|
||||
EventsGuard events_{2};
|
||||
hipError_t error_;
|
||||
|
||||
char* out_ = nullptr;
|
||||
void* small_kernel_args_[2] = {&small_kernel_args, &out_};
|
||||
void* medium_kernel_args_[2] = {&medium_kernel_args, &out_};
|
||||
void* large_kernel_args_[2] = {&large_kernel_args, &out_};
|
||||
};
|
||||
|
||||
template <KernelType kernel_type, bool timer_type> static void RunBenchmark(bool sync) {
|
||||
ExtLaunchKernelBenchmark<kernel_type, timer_type> benchmark;
|
||||
benchmark.AddSectionName(GetSynchronizationSectionName(sync));
|
||||
benchmark.AddSectionName(GetKernelTypeSectionName<kernel_type>());
|
||||
benchmark.AddSectionName(GetTimerTypeSectionName<timer_type>());
|
||||
benchmark.Run(sync);
|
||||
HIP_CHECK(benchmark.GetError());
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Calls an empty kernel using hipExtLaunchKernel:
|
||||
* -# With different timing methods:
|
||||
* - CPU-based
|
||||
* - Event-based
|
||||
* -# With different synchronization behavior:
|
||||
* - Using a stream synchronization between each iteration
|
||||
* - Without any synchronization between iterations
|
||||
* -# With different kernel argument sizes
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/kernelLaunch/hipExtLaunchKernel.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipExtLaunchKernel") {
|
||||
bool sync = GENERATE(true, false);
|
||||
|
||||
SECTION("null kernel") {
|
||||
SECTION("cpu-based timing") { RunBenchmark<KernelType::kNull, kTimerTypeCpu>(sync); }
|
||||
|
||||
SECTION("event-based timing") { RunBenchmark<KernelType::kNull, kTimerTypeEvent>(sync); }
|
||||
}
|
||||
|
||||
SECTION("small kernel") {
|
||||
SECTION("cpu-based timing") { RunBenchmark<KernelType::kSmall, kTimerTypeCpu>(sync); }
|
||||
|
||||
SECTION("event-based timing") { RunBenchmark<KernelType::kSmall, kTimerTypeEvent>(sync); }
|
||||
}
|
||||
|
||||
SECTION("medium kernel") {
|
||||
SECTION("cpu-based timing") { RunBenchmark<KernelType::kMedium, kTimerTypeCpu>(sync); }
|
||||
|
||||
SECTION("event-based timing") { RunBenchmark<KernelType::kMedium, kTimerTypeEvent>(sync); }
|
||||
}
|
||||
|
||||
SECTION("large kernel") {
|
||||
SECTION("cpu-based timing") { RunBenchmark<KernelType::kLarge, kTimerTypeCpu>(sync); }
|
||||
|
||||
SECTION("event-based timing") { RunBenchmark<KernelType::kLarge, kTimerTypeEvent>(sync); }
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,130 @@
|
||||
/*
|
||||
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 "kernel_launch_common.hh"
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <utils.hh>
|
||||
|
||||
/**
|
||||
* @addtogroup kernelLaunch kernel launch
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
* Contains performance tests for kernel launch overhead benchmarking.
|
||||
*/
|
||||
|
||||
template <KernelType kernel_type, bool timer_type>
|
||||
class LaunchCooperativeKernelBenchmark
|
||||
: public KernelLaunchBenchmark<LaunchCooperativeKernelBenchmark<kernel_type, timer_type>,
|
||||
timer_type> {
|
||||
public:
|
||||
constexpr void LaunchKernel() {
|
||||
if constexpr (kernel_type == KernelType::kNull) {
|
||||
error_ = hipLaunchCooperativeKernel(reinterpret_cast<void*>(NullKernel), dim3{1, 1, 1},
|
||||
dim3{1, 1, 1}, nullptr, 0, nullptr);
|
||||
} else if constexpr (kernel_type == KernelType::kSmall) {
|
||||
error_ =
|
||||
hipLaunchCooperativeKernel(reinterpret_cast<void*>(KernelWithSmallArgs), dim3{1, 1, 1},
|
||||
dim3{1, 1, 1}, small_kernel_args_, 0, nullptr);
|
||||
} else if constexpr (kernel_type == KernelType::kMedium) {
|
||||
error_ =
|
||||
hipLaunchCooperativeKernel(reinterpret_cast<void*>(KernelWithMediumArgs), dim3{1, 1, 1},
|
||||
dim3{1, 1, 1}, medium_kernel_args_, 0, nullptr);
|
||||
} else if constexpr (kernel_type == KernelType::kLarge) {
|
||||
error_ =
|
||||
hipLaunchCooperativeKernel(reinterpret_cast<void*>(KernelWithLargeArgs), dim3{1, 1, 1},
|
||||
dim3{1, 1, 1}, large_kernel_args_, 0, nullptr);
|
||||
} else
|
||||
;
|
||||
}
|
||||
|
||||
hipError_t GetError() { return error_; }
|
||||
|
||||
private:
|
||||
hipError_t error_;
|
||||
|
||||
char* out_ = nullptr;
|
||||
void* small_kernel_args_[2] = {&small_kernel_args, &out_};
|
||||
void* medium_kernel_args_[2] = {&medium_kernel_args, &out_};
|
||||
void* large_kernel_args_[2] = {&large_kernel_args, &out_};
|
||||
};
|
||||
|
||||
template <KernelType kernel_type, bool timer_type> static void RunBenchmark(bool sync) {
|
||||
LaunchCooperativeKernelBenchmark<kernel_type, timer_type> benchmark;
|
||||
benchmark.AddSectionName(GetSynchronizationSectionName(sync));
|
||||
benchmark.AddSectionName(GetKernelTypeSectionName<kernel_type>());
|
||||
benchmark.AddSectionName(GetTimerTypeSectionName<timer_type>());
|
||||
benchmark.Run(sync);
|
||||
HIP_CHECK(benchmark.GetError());
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Calls an empty kernel using hipLaunchCooperativeKernel:
|
||||
* -# With different timing methods:
|
||||
* - CPU-based
|
||||
* - Event-based
|
||||
* -# With different synchronization behavior:
|
||||
* - Using a stream synchronization between each iteration
|
||||
* - Without any synchronization between iterations
|
||||
* -# With different kernel argument sizes
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/kernelLaunch/hipLaunchCooperativeKernel.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Device supports CooperativeLaunch
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipLaunchCooperativeKernel") {
|
||||
if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) {
|
||||
HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported");
|
||||
return;
|
||||
}
|
||||
|
||||
bool sync = GENERATE(true, false);
|
||||
|
||||
SECTION("null kernel") {
|
||||
SECTION("cpu-based timing") { RunBenchmark<KernelType::kNull, kTimerTypeCpu>(sync); }
|
||||
|
||||
SECTION("event-based timing") { RunBenchmark<KernelType::kNull, kTimerTypeEvent>(sync); }
|
||||
}
|
||||
|
||||
SECTION("small kernel") {
|
||||
SECTION("cpu-based timing") { RunBenchmark<KernelType::kSmall, kTimerTypeCpu>(sync); }
|
||||
|
||||
SECTION("event-based timing") { RunBenchmark<KernelType::kSmall, kTimerTypeEvent>(sync); }
|
||||
}
|
||||
|
||||
SECTION("medium kernel") {
|
||||
SECTION("cpu-based timing") { RunBenchmark<KernelType::kMedium, kTimerTypeCpu>(sync); }
|
||||
|
||||
SECTION("event-based timing") { RunBenchmark<KernelType::kMedium, kTimerTypeEvent>(sync); }
|
||||
}
|
||||
|
||||
SECTION("large kernel") {
|
||||
SECTION("cpu-based timing") { RunBenchmark<KernelType::kLarge, kTimerTypeCpu>(sync); }
|
||||
|
||||
SECTION("event-based timing") { RunBenchmark<KernelType::kLarge, kTimerTypeEvent>(sync); }
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,118 @@
|
||||
/*
|
||||
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 "kernel_launch_common.hh"
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
/**
|
||||
* @addtogroup kernelLaunch kernel launch
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
* Contains performance tests for kernel launch overhead benchmarking.
|
||||
*/
|
||||
|
||||
template <KernelType kernel_type, bool timer_type>
|
||||
class LaunchKernelBenchmark
|
||||
: public KernelLaunchBenchmark<LaunchKernelBenchmark<kernel_type, timer_type>, timer_type> {
|
||||
public:
|
||||
constexpr void LaunchKernel() {
|
||||
if constexpr (kernel_type == KernelType::kNull) {
|
||||
error_ = hipLaunchKernel(reinterpret_cast<void*>(NullKernel), 1, 1, nullptr, 0, nullptr);
|
||||
} else if constexpr (kernel_type == KernelType::kSmall) {
|
||||
error_ = hipLaunchKernel(reinterpret_cast<void*>(KernelWithSmallArgs), 1, 1,
|
||||
small_kernel_args_, 0, nullptr);
|
||||
} else if constexpr (kernel_type == KernelType::kMedium) {
|
||||
error_ = hipLaunchKernel(reinterpret_cast<void*>(KernelWithMediumArgs), 1, 1,
|
||||
medium_kernel_args_, 0, nullptr);
|
||||
} else if constexpr (kernel_type == KernelType::kLarge) {
|
||||
error_ = hipLaunchKernel(reinterpret_cast<void*>(KernelWithLargeArgs), 1, 1,
|
||||
large_kernel_args_, 0, nullptr);
|
||||
} else
|
||||
;
|
||||
}
|
||||
|
||||
hipError_t GetError() { return error_; }
|
||||
|
||||
private:
|
||||
hipError_t error_;
|
||||
|
||||
char* out_ = nullptr;
|
||||
void* small_kernel_args_[2] = {&small_kernel_args, &out_};
|
||||
void* medium_kernel_args_[2] = {&medium_kernel_args, &out_};
|
||||
void* large_kernel_args_[2] = {&large_kernel_args, &out_};
|
||||
};
|
||||
|
||||
template <KernelType kernel_type, bool timer_type> static void RunBenchmark(bool sync) {
|
||||
LaunchKernelBenchmark<kernel_type, timer_type> benchmark;
|
||||
benchmark.AddSectionName(GetSynchronizationSectionName(sync));
|
||||
benchmark.AddSectionName(GetKernelTypeSectionName<kernel_type>());
|
||||
benchmark.AddSectionName(GetTimerTypeSectionName<timer_type>());
|
||||
benchmark.Run(sync);
|
||||
HIP_CHECK(benchmark.GetError());
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Calls an empty kernel using hipLaunchKernel:
|
||||
* -# With different timing methods:
|
||||
* - CPU-based
|
||||
* - Event-based
|
||||
* -# With different synchronization behavior:
|
||||
* - Using a stream synchronization between each iteration
|
||||
* - Without any synchronization between iterations
|
||||
* -# With different kernel argument sizes
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/kernelLaunch/hipLaunchKernel.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipLaunchKernel") {
|
||||
bool sync = GENERATE(true, false);
|
||||
|
||||
SECTION("null kernel") {
|
||||
SECTION("cpu-based timing") { RunBenchmark<KernelType::kNull, kTimerTypeCpu>(sync); }
|
||||
|
||||
SECTION("event-based timing") { RunBenchmark<KernelType::kNull, kTimerTypeEvent>(sync); }
|
||||
}
|
||||
|
||||
SECTION("small kernel") {
|
||||
SECTION("cpu-based timing") { RunBenchmark<KernelType::kSmall, kTimerTypeCpu>(sync); }
|
||||
|
||||
SECTION("event-based timing") { RunBenchmark<KernelType::kSmall, kTimerTypeEvent>(sync); }
|
||||
}
|
||||
|
||||
SECTION("medium kernel") {
|
||||
SECTION("cpu-based timing") { RunBenchmark<KernelType::kMedium, kTimerTypeCpu>(sync); }
|
||||
|
||||
SECTION("event-based timing") { RunBenchmark<KernelType::kMedium, kTimerTypeEvent>(sync); }
|
||||
}
|
||||
|
||||
SECTION("large kernel") {
|
||||
SECTION("cpu-based timing") { RunBenchmark<KernelType::kLarge, kTimerTypeCpu>(sync); }
|
||||
|
||||
SECTION("event-based timing") { RunBenchmark<KernelType::kLarge, kTimerTypeEvent>(sync); }
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,39 @@
|
||||
/*
|
||||
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 "kernel_launch_common.hh"
|
||||
|
||||
#define DO_NOT_OPTIMIZE_AWAY \
|
||||
unsigned i = blockIdx.x * blockDim.x + threadIdx.x; \
|
||||
if (out) *out = args.args[i];
|
||||
|
||||
__global__ void NullKernel() {}
|
||||
|
||||
__global__ void KernelWithSmallArgs(SmallKernelArgs args, char* out) { DO_NOT_OPTIMIZE_AWAY; }
|
||||
|
||||
__global__ void KernelWithMediumArgs(MediumKernelArgs args, char* out) { DO_NOT_OPTIMIZE_AWAY; }
|
||||
|
||||
__global__ void KernelWithLargeArgs(LargeKernelArgs args, char* out) { DO_NOT_OPTIMIZE_AWAY; }
|
||||
|
||||
SmallKernelArgs small_kernel_args;
|
||||
MediumKernelArgs medium_kernel_args;
|
||||
LargeKernelArgs large_kernel_args;
|
||||
@@ -0,0 +1,116 @@
|
||||
/*
|
||||
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
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <performance_common.hh>
|
||||
|
||||
struct SmallKernelArgs {
|
||||
char args[16];
|
||||
};
|
||||
|
||||
struct MediumKernelArgs {
|
||||
char args[256];
|
||||
};
|
||||
|
||||
struct LargeKernelArgs {
|
||||
char args[4080];
|
||||
};
|
||||
|
||||
extern SmallKernelArgs small_kernel_args;
|
||||
extern MediumKernelArgs medium_kernel_args;
|
||||
extern LargeKernelArgs large_kernel_args;
|
||||
|
||||
__global__ void NullKernel();
|
||||
|
||||
__global__ void KernelWithSmallArgs(SmallKernelArgs, char*);
|
||||
|
||||
__global__ void KernelWithMediumArgs(MediumKernelArgs, char*);
|
||||
|
||||
__global__ void KernelWithLargeArgs(LargeKernelArgs, char*);
|
||||
|
||||
enum class KernelType { kNull = 0, kSmall, kMedium, kLarge };
|
||||
|
||||
template <typename Derived, bool timer_type>
|
||||
class KernelLaunchBenchmark : public Benchmark<KernelLaunchBenchmark<Derived, timer_type>> {
|
||||
public:
|
||||
void operator()(bool sync = true) {
|
||||
auto& derived = static_cast<Derived&>(*this);
|
||||
|
||||
if (sync) {
|
||||
TIMED_SECTION(timer_type) { derived.LaunchKernel(); }
|
||||
} else {
|
||||
if (this->current() != this->kWarmup) // if not warmup
|
||||
RunWithoutSynchronization();
|
||||
}
|
||||
}
|
||||
|
||||
private:
|
||||
void RunWithoutSynchronization() {
|
||||
auto iterations = this->iterations();
|
||||
auto warmups = this->warmups();
|
||||
|
||||
// manually handle iterations here to avoid synchronization after each iteration
|
||||
this->Configure(1, 0);
|
||||
|
||||
this->RegisterModifier([iterations](float time) { return time / iterations; });
|
||||
|
||||
auto& derived = static_cast<Derived&>(*this);
|
||||
|
||||
for (size_t i = 0u; i < warmups; ++i) {
|
||||
derived.LaunchKernel();
|
||||
}
|
||||
|
||||
TIMED_SECTION(timer_type) {
|
||||
for (size_t i = 0u; i < iterations; ++i) {
|
||||
derived.LaunchKernel();
|
||||
}
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
static std::string GetSynchronizationSectionName(bool sync) {
|
||||
return sync ? "with synchronization" : "without synchronization";
|
||||
}
|
||||
|
||||
template <KernelType kernel_type> std::string GetKernelTypeSectionName() {
|
||||
if constexpr (kernel_type == KernelType::kNull) {
|
||||
return "null kernel";
|
||||
} else if constexpr (kernel_type == KernelType::kSmall) {
|
||||
return "small kernel";
|
||||
} else if constexpr (kernel_type == KernelType::kMedium) {
|
||||
return "medium kernel";
|
||||
} else if constexpr (kernel_type == KernelType::kLarge) {
|
||||
return "large kernel";
|
||||
} else {
|
||||
return "unknown kernel type";
|
||||
}
|
||||
}
|
||||
|
||||
template <bool timer_type> std::string GetTimerTypeSectionName() {
|
||||
if constexpr (timer_type == kTimerTypeEvent) {
|
||||
return "event based";
|
||||
} else {
|
||||
return "cpu based";
|
||||
}
|
||||
}
|
||||
@@ -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 "kernel_launch_common.hh"
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
/**
|
||||
* @addtogroup kernelLaunch kernel launch
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
* Contains performance tests for kernel launch overhead benchmarking.
|
||||
*/
|
||||
|
||||
template <KernelType kernel_type, bool timer_type>
|
||||
class TripleChevronBenchmark
|
||||
: public KernelLaunchBenchmark<TripleChevronBenchmark<kernel_type, timer_type>, timer_type> {
|
||||
public:
|
||||
constexpr void LaunchKernel() {
|
||||
if constexpr (kernel_type == KernelType::kNull) {
|
||||
NullKernel<<<1, 1>>>();
|
||||
} else if constexpr (kernel_type == KernelType::kSmall) {
|
||||
KernelWithSmallArgs<<<1, 1>>>(small_kernel_args, nullptr);
|
||||
} else if constexpr (kernel_type == KernelType::kMedium) {
|
||||
KernelWithMediumArgs<<<1, 1>>>(medium_kernel_args, nullptr);
|
||||
} else if constexpr (kernel_type == KernelType::kLarge) {
|
||||
KernelWithLargeArgs<<<1, 1>>>(large_kernel_args, nullptr);
|
||||
} else
|
||||
;
|
||||
}
|
||||
};
|
||||
|
||||
template <KernelType kernel_type, bool timer_type> static void RunBenchmark(bool sync) {
|
||||
TripleChevronBenchmark<kernel_type, timer_type> benchmark;
|
||||
benchmark.AddSectionName(GetSynchronizationSectionName(sync));
|
||||
benchmark.AddSectionName(GetKernelTypeSectionName<kernel_type>());
|
||||
benchmark.AddSectionName(GetTimerTypeSectionName<timer_type>());
|
||||
benchmark.Run(sync);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Calls an empty kernel using triple chevron annotation:
|
||||
* -# With different timing methods:
|
||||
* - CPU-based
|
||||
* - Event-based
|
||||
* -# With different synchronization behavior:
|
||||
* - Using a stream synchronization between each iteration
|
||||
* - Without any synchronization between iterations
|
||||
* -# With different kernel argument sizes
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/kernelLaunch/triple_chevron.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_Triple_Chevron") {
|
||||
bool sync = GENERATE(true, false);
|
||||
|
||||
SECTION("null kernel") {
|
||||
SECTION("cpu-based timing") { RunBenchmark<KernelType::kNull, kTimerTypeCpu>(sync); }
|
||||
|
||||
SECTION("event-based timing") { RunBenchmark<KernelType::kNull, kTimerTypeEvent>(sync); }
|
||||
}
|
||||
|
||||
SECTION("small kernel") {
|
||||
SECTION("cpu-based timing") { RunBenchmark<KernelType::kSmall, kTimerTypeCpu>(sync); }
|
||||
|
||||
SECTION("event-based timing") { RunBenchmark<KernelType::kSmall, kTimerTypeEvent>(sync); }
|
||||
}
|
||||
|
||||
SECTION("medium kernel") {
|
||||
SECTION("cpu-based timing") { RunBenchmark<KernelType::kMedium, kTimerTypeCpu>(sync); }
|
||||
|
||||
SECTION("event-based timing") { RunBenchmark<KernelType::kMedium, kTimerTypeEvent>(sync); }
|
||||
}
|
||||
|
||||
SECTION("large kernel") {
|
||||
SECTION("cpu-based timing") { RunBenchmark<KernelType::kLarge, kTimerTypeCpu>(sync); }
|
||||
|
||||
SECTION("event-based timing") { RunBenchmark<KernelType::kLarge, kTimerTypeEvent>(sync); }
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,52 @@
|
||||
# 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.
|
||||
|
||||
set(TEST_SRC
|
||||
hipMemcpy.cc
|
||||
hipMemcpyAsync.cc
|
||||
hipMemcpyWithStream.cc
|
||||
hipMemcpyAtoH.cc
|
||||
hipMemcpyHtoA.cc
|
||||
hipMemcpyDtoD.cc
|
||||
hipMemcpyDtoDAsync.cc
|
||||
hipMemcpyDtoH.cc
|
||||
hipMemcpyDtoHAsync.cc
|
||||
hipMemcpyHtoD.cc
|
||||
hipMemcpyHtoDAsync.cc
|
||||
hipMemcpyToSymbol.cc
|
||||
hipMemcpyToSymbolAsync.cc
|
||||
hipMemcpyFromSymbol.cc
|
||||
hipMemcpyFromSymbolAsync.cc
|
||||
hipMemcpy2D.cc
|
||||
hipMemcpy2DAsync.cc
|
||||
hipMemcpy2DToArray.cc
|
||||
hipMemcpy2DToArrayAsync.cc
|
||||
hipMemcpy2DFromArray.cc
|
||||
hipMemcpy2DFromArrayAsync.cc
|
||||
hipMemcpyParam2D.cc
|
||||
hipMemcpyParam2DAsync.cc
|
||||
hipMemcpy3D.cc
|
||||
hipMemcpy3DAsync.cc
|
||||
)
|
||||
|
||||
hip_add_exe_to_target(NAME MemcpyPerformance
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME build_tests
|
||||
COMPILE_OPTIONS -std=c++17)
|
||||
@@ -0,0 +1,190 @@
|
||||
/*
|
||||
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 "memcpy_performance_common.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
* Contains performance tests for all memcpy HIP APIs.
|
||||
*/
|
||||
|
||||
class MemcpyBenchmark : public Benchmark<MemcpyBenchmark> {
|
||||
public:
|
||||
void operator()(void* dst, const void* src, size_t size, hipMemcpyKind kind) {
|
||||
TIMED_SECTION(kTimerTypeCpu) {
|
||||
HIP_CHECK(hipMemcpy(dst, src, size, kind));
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(LinearAllocs dst_allocation_type, LinearAllocs src_allocation_type,
|
||||
size_t size, hipMemcpyKind kind, bool enable_peer_access=false) {
|
||||
MemcpyBenchmark benchmark;
|
||||
benchmark.AddSectionName(std::to_string(size));
|
||||
benchmark.AddSectionName(GetAllocationSectionName(src_allocation_type));
|
||||
benchmark.AddSectionName(GetAllocationSectionName(dst_allocation_type));
|
||||
|
||||
if (kind != hipMemcpyDeviceToDevice) {
|
||||
LinearAllocGuard<int> src_allocation(src_allocation_type, size);
|
||||
LinearAllocGuard<int> dst_allocation(dst_allocation_type, size);
|
||||
benchmark.Run(dst_allocation.ptr(), src_allocation.ptr(), size, kind);
|
||||
} else {
|
||||
int src_device = std::get<0>(GetDeviceIds(enable_peer_access));
|
||||
int dst_device = std::get<1>(GetDeviceIds(enable_peer_access));
|
||||
|
||||
LinearAllocGuard<int> src_allocation(src_allocation_type, size);
|
||||
HIP_CHECK(hipSetDevice(dst_device));
|
||||
LinearAllocGuard<int> dst_allocation(dst_allocation_type, size);
|
||||
HIP_CHECK(hipSetDevice(src_device));
|
||||
benchmark.Run(dst_allocation.ptr(), src_allocation.ptr(), size, kind);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy` from Device to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: device malloc
|
||||
* - Destination: host pinned and pageable
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy_DeviceToHost") {
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto src_allocation_type = LinearAllocs::hipMalloc;
|
||||
const auto dst_allocation_type = GENERATE(LinearAllocs::malloc, LinearAllocs::hipHostMalloc);
|
||||
RunBenchmark(dst_allocation_type, src_allocation_type, allocation_size, hipMemcpyDeviceToHost);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy` from Host to Device:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: host pinned and pageable
|
||||
* - Destination: device malloc
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy_HostToDevice") {
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto src_allocation_type = GENERATE(LinearAllocs::malloc, LinearAllocs::hipHostMalloc);
|
||||
const auto dst_allocation_type = LinearAllocs::hipMalloc;
|
||||
RunBenchmark(dst_allocation_type, src_allocation_type, allocation_size, hipMemcpyHostToDevice);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy` from Host to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: host pinned and pageable
|
||||
* - Destination: host pinned and pageable
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy_HostToHost") {
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto src_allocation_type = GENERATE(LinearAllocs::malloc, LinearAllocs::hipHostMalloc);
|
||||
const auto dst_allocation_type = GENERATE(LinearAllocs::malloc, LinearAllocs::hipHostMalloc);
|
||||
RunBenchmark(dst_allocation_type, src_allocation_type, allocation_size, hipMemcpyHostToHost);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy` from Device to Device with peer access enabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: device malloc
|
||||
* - Destination: device malloc
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Multi-device
|
||||
* - Device supports Peer-to-Peer access
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy_DeviceToDevice_EnablePeerAccess") {
|
||||
if (HipTest::getDeviceCount() < 2) {
|
||||
HipTest::HIP_SKIP_TEST("This test requires 2 GPUs. Skipping.");
|
||||
return;
|
||||
}
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto src_allocation_type = LinearAllocs::hipMalloc;
|
||||
const auto dst_allocation_type = LinearAllocs::hipMalloc;
|
||||
RunBenchmark(dst_allocation_type, src_allocation_type, allocation_size, hipMemcpyDeviceToDevice, true);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy` from Device to Device with peer access disabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: device malloc
|
||||
* - Destination: device malloc
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy_DeviceToDevice_DisablePeerAccess") {
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto src_allocation_type = LinearAllocs::hipMalloc;
|
||||
const auto dst_allocation_type = LinearAllocs::hipMalloc;
|
||||
RunBenchmark(dst_allocation_type, src_allocation_type, allocation_size, hipMemcpyDeviceToDevice);
|
||||
}
|
||||
@@ -0,0 +1,183 @@
|
||||
/*
|
||||
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 "memcpy_performance_common.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class Memcpy2DBenchmark : public Benchmark<Memcpy2DBenchmark> {
|
||||
public:
|
||||
void operator()(void* dst, size_t dst_pitch, const void* src, size_t src_pitch, size_t width,
|
||||
size_t height, hipMemcpyKind kind) {
|
||||
TIMED_SECTION(kTimerTypeCpu) {
|
||||
HIP_CHECK(hipMemcpy2D(dst, dst_pitch, src, src_pitch, width, height, kind));
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(size_t width, size_t height, hipMemcpyKind kind, bool enable_peer_access=false) {
|
||||
Memcpy2DBenchmark benchmark;
|
||||
benchmark.AddSectionName("(" + std::to_string(width) + ", " + std::to_string(height) + ")");
|
||||
|
||||
if (kind == hipMemcpyDeviceToHost) {
|
||||
LinearAllocGuard2D<int> device_allocation(width, height);
|
||||
LinearAllocGuard<int> host_allocation(LinearAllocs::hipHostMalloc,
|
||||
device_allocation.width() * height);
|
||||
benchmark.Run(host_allocation.ptr(), device_allocation.width(),
|
||||
device_allocation.ptr(), device_allocation.pitch(),
|
||||
device_allocation.width(), device_allocation.height(),
|
||||
hipMemcpyDeviceToHost);
|
||||
} else if (kind == hipMemcpyHostToDevice) {
|
||||
LinearAllocGuard2D<int> device_allocation(width, height);
|
||||
LinearAllocGuard<int> host_allocation(LinearAllocs::hipHostMalloc,
|
||||
device_allocation.width() * height);
|
||||
benchmark.Run(device_allocation.ptr(), device_allocation.pitch(),
|
||||
host_allocation.ptr(), device_allocation.width(),
|
||||
device_allocation.width(), device_allocation.height(),
|
||||
hipMemcpyHostToDevice);
|
||||
} else if (kind == hipMemcpyHostToHost) {
|
||||
LinearAllocGuard<int> src_allocation(LinearAllocs::hipHostMalloc, width * sizeof(int) * height);
|
||||
LinearAllocGuard<int> dst_allocation(LinearAllocs::hipHostMalloc, width * sizeof(int) * height);
|
||||
benchmark.Run(dst_allocation.ptr(), width * sizeof(int), src_allocation.ptr(),
|
||||
width * sizeof(int), width * sizeof(int), height, hipMemcpyHostToHost);
|
||||
} else {
|
||||
// hipMemcpyDeviceToDevice
|
||||
int src_device = std::get<0>(GetDeviceIds(enable_peer_access));
|
||||
int dst_device = std::get<1>(GetDeviceIds(enable_peer_access));
|
||||
|
||||
LinearAllocGuard2D<int> src_allocation(width, height);
|
||||
HIP_CHECK(hipSetDevice(dst_device));
|
||||
LinearAllocGuard2D<int> dst_allocation(width, height);
|
||||
HIP_CHECK(hipSetDevice(src_device));
|
||||
benchmark.Run(dst_allocation.ptr(), dst_allocation.pitch(),
|
||||
src_allocation.ptr(), src_allocation.pitch(),
|
||||
dst_allocation.width(), dst_allocation.height(),
|
||||
hipMemcpyDeviceToDevice);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy2D` from Device to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 4 MB x 32 B
|
||||
* - Large: 16 MB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy2D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy2D_DeviceToHost") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 32, hipMemcpyDeviceToHost);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy2D` from Host to Device:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 4 MB x 32 B
|
||||
* - Large: 16 MB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy2D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy2D_HostToDevice") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 32, hipMemcpyHostToDevice);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy2D` from Host to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 4 MB x 32 B
|
||||
* - Large: 16 MB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy2D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy2D_HostToHost") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 32, hipMemcpyHostToHost);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy2D` from Device to Device with peer access disabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 4 MB x 32 B
|
||||
* - Large: 16 MB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy2D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy2D_DeviceToDevice_DisablePeerAccess") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 32, hipMemcpyDeviceToDevice);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy2D` from Device to Device with peer access enabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 4 MB x 32 B
|
||||
* - Large: 16 MB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy2D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Multi-device
|
||||
* - Device supports Peer-to-Peer access
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy2D_DeviceToDevice_EnablePeerAccess") {
|
||||
if (HipTest::getDeviceCount() < 2) {
|
||||
HipTest::HIP_SKIP_TEST("This test requires 2 GPUs. Skipping.");
|
||||
return;
|
||||
}
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 32, hipMemcpyDeviceToDevice, true);
|
||||
}
|
||||
@@ -0,0 +1,188 @@
|
||||
/*
|
||||
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 "memcpy_performance_common.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class Memcpy2DAsyncBenchmark : public Benchmark<Memcpy2DAsyncBenchmark> {
|
||||
public:
|
||||
void operator()(void* dst, size_t dst_pitch, const void* src, size_t src_pitch, size_t width,
|
||||
size_t height, hipMemcpyKind kind, const hipStream_t& stream) {
|
||||
TIMED_SECTION_STREAM(kTimerTypeEvent, stream) {
|
||||
HIP_CHECK(hipMemcpy2DAsync(dst, dst_pitch, src, src_pitch, width, height, kind, stream));
|
||||
}
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(size_t width, size_t height, hipMemcpyKind kind, bool enable_peer_access=false) {
|
||||
Memcpy2DAsyncBenchmark benchmark;
|
||||
benchmark.AddSectionName("(" + std::to_string(width) + ", " + std::to_string(height) + ")");
|
||||
|
||||
const StreamGuard stream_guard(Streams::created);
|
||||
const hipStream_t stream = stream_guard.stream();
|
||||
|
||||
if (kind == hipMemcpyDeviceToHost) {
|
||||
LinearAllocGuard2D<int> device_allocation(width, height);
|
||||
LinearAllocGuard<int> host_allocation(LinearAllocs::hipHostMalloc,
|
||||
device_allocation.width() * height);
|
||||
benchmark.Run(host_allocation.ptr(), device_allocation.width(),
|
||||
device_allocation.ptr(), device_allocation.pitch(),
|
||||
device_allocation.width(), device_allocation.height(),
|
||||
hipMemcpyDeviceToHost, stream);
|
||||
} else if (kind == hipMemcpyHostToDevice) {
|
||||
LinearAllocGuard2D<int> device_allocation(width, height);
|
||||
LinearAllocGuard<int> host_allocation(LinearAllocs::hipHostMalloc,
|
||||
device_allocation.width() * height);
|
||||
benchmark.Run(device_allocation.ptr(), device_allocation.pitch(),
|
||||
host_allocation.ptr(), device_allocation.width(),
|
||||
device_allocation.width(), device_allocation.height(),
|
||||
hipMemcpyHostToDevice, stream);
|
||||
} else if (kind == hipMemcpyHostToHost) {
|
||||
LinearAllocGuard<int> src_allocation(LinearAllocs::hipHostMalloc, width * sizeof(int) * height);
|
||||
LinearAllocGuard<int> dst_allocation(LinearAllocs::hipHostMalloc, width * sizeof(int) * height);
|
||||
benchmark.Run(dst_allocation.ptr(), width * sizeof(int), src_allocation.ptr(),
|
||||
width * sizeof(int), width * sizeof(int), height, hipMemcpyHostToHost, stream);
|
||||
} else {
|
||||
// hipMemcpyDeviceToDevice
|
||||
int src_device = std::get<0>(GetDeviceIds(enable_peer_access));
|
||||
int dst_device = std::get<1>(GetDeviceIds(enable_peer_access));
|
||||
|
||||
LinearAllocGuard2D<int> src_allocation(width, height);
|
||||
HIP_CHECK(hipSetDevice(dst_device));
|
||||
LinearAllocGuard2D<int> dst_allocation(width, height);
|
||||
|
||||
HIP_CHECK(hipSetDevice(src_device));
|
||||
benchmark.Run(dst_allocation.ptr(), dst_allocation.pitch(),
|
||||
src_allocation.ptr(), src_allocation.pitch(),
|
||||
dst_allocation.width(), dst_allocation.height(),
|
||||
hipMemcpyDeviceToDevice, stream);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy2DAsync` from Device to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 4 MB x 32 B
|
||||
* - Large: 16 MB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy2DAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy2DAsync_DeviceToHost") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 32, hipMemcpyDeviceToHost);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy2DAsync` from Host to Device:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 4 MB x 32 B
|
||||
* - Large: 16 MB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy2DAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy2DAsync_HostToDevice") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 32, hipMemcpyHostToDevice);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy2DAsync` from Host to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 4 MB x 32 B
|
||||
* - Large: 16 MB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy2DAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy2DAsync_HostToHost") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 32, hipMemcpyHostToHost);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy2DAsync` from Device to Device with peer access disabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 4 MB x 32 B
|
||||
* - Large: 16 MB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy2DAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy2DAsync_DeviceToDevice_DisablePeerAccess") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 32, hipMemcpyDeviceToDevice);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy2D` from Device to Device with peer access enabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 4 MB x 32 B
|
||||
* - Large: 16 MB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy2DAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Multi-device
|
||||
* - Device supports Peer-to-Peer access
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy2DAsync_DeviceToDevice_EnablePeerAccess") {
|
||||
if (HipTest::getDeviceCount() < 2) {
|
||||
HipTest::HIP_SKIP_TEST("This test requires 2 GPUs. Skipping.");
|
||||
return;
|
||||
}
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 32, hipMemcpyDeviceToDevice, true);
|
||||
}
|
||||
@@ -0,0 +1,127 @@
|
||||
/*
|
||||
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 "memcpy_performance_common.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class Memcpy2DFromArrayBenchmark : public Benchmark<Memcpy2DFromArrayBenchmark> {
|
||||
public:
|
||||
void operator()(void* dst, size_t dst_pitch, hipArray_const_t src, size_t width, size_t height, hipMemcpyKind kind) {
|
||||
TIMED_SECTION(kTimerTypeCpu) {
|
||||
HIP_CHECK(hipMemcpy2DFromArray(dst, dst_pitch, src, 0, 0, width, height, kind));
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(size_t width, size_t height, hipMemcpyKind kind,
|
||||
bool enable_peer_access=false) {
|
||||
Memcpy2DFromArrayBenchmark benchmark;
|
||||
benchmark.AddSectionName("(" + std::to_string(width) + ", " + std::to_string(height) + ")");
|
||||
|
||||
if (kind == hipMemcpyDeviceToHost) {
|
||||
size_t allocation_size = width * height * sizeof(int);
|
||||
LinearAllocGuard<int> host_allocation(LinearAllocs::hipHostMalloc, allocation_size);
|
||||
ArrayAllocGuard<int> array_allocation(make_hipExtent(width, height, 0), hipArrayDefault);
|
||||
benchmark.Run(host_allocation.ptr(), width * sizeof(int), array_allocation.ptr(),
|
||||
width * sizeof(int), height, hipMemcpyDeviceToHost);
|
||||
} else {
|
||||
// hipMemcpyDeviceToDevice
|
||||
int src_device = std::get<0>(GetDeviceIds(enable_peer_access));
|
||||
int dst_device = std::get<1>(GetDeviceIds(enable_peer_access));
|
||||
|
||||
LinearAllocGuard2D<int> device_allocation(width, height);
|
||||
HIP_CHECK(hipSetDevice(dst_device));
|
||||
ArrayAllocGuard<int> array_allocation(make_hipExtent(width, height, 0), hipArrayDefault);
|
||||
HIP_CHECK(hipSetDevice(src_device));
|
||||
benchmark.Run(device_allocation.ptr(), device_allocation.pitch(),
|
||||
array_allocation.ptr(), device_allocation.width(),
|
||||
device_allocation.height(), hipMemcpyDeviceToDevice);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy2DFromArray` from Device to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 8 KB x 32 B
|
||||
* - Large: 16 KB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy2DFromArray.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy2DFromArray_DeviceToHost") {
|
||||
const auto width = GENERATE(4_KB, 8_KB, 16_KB);
|
||||
RunBenchmark(width, 32, hipMemcpyDeviceToHost);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy2DFromArray` from Device to Device with peer access disabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 8 KB x 32 B
|
||||
* - Large: 16 KB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy2DFromArray.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy2DFromArray_DeviceToDevice_DisablePeerAccess") {
|
||||
const auto width = GENERATE(4_KB, 8_KB, 16_KB);
|
||||
RunBenchmark(width, 32, hipMemcpyDeviceToDevice);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy2DFromArray` from Device to Device with peer access enabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 8 KB x 32 B
|
||||
* - Large: 16 KB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy2DFromArray.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Multi-device
|
||||
* - Device supports Peer-to-Peer access
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy2DFromArray_DeviceToDevice_EnablePeerAccess") {
|
||||
if (HipTest::getDeviceCount() < 2) {
|
||||
HipTest::HIP_SKIP_TEST("This test requires 2 GPUs. Skipping.");
|
||||
return;
|
||||
}
|
||||
const auto width = GENERATE(4_KB, 8_KB, 16_KB);
|
||||
RunBenchmark(width, 32, hipMemcpyDeviceToDevice, true);
|
||||
}
|
||||
@@ -0,0 +1,133 @@
|
||||
/*
|
||||
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 "memcpy_performance_common.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class Memcpy2DFromArrayAsyncBenchmark : public Benchmark<Memcpy2DFromArrayAsyncBenchmark> {
|
||||
public:
|
||||
void operator()(void* dst, size_t dst_pitch, hipArray_const_t src, size_t width, size_t height,
|
||||
hipMemcpyKind kind, const hipStream_t& stream) {
|
||||
TIMED_SECTION_STREAM(kTimerTypeEvent, stream) {
|
||||
HIP_CHECK(hipMemcpy2DFromArrayAsync(dst, dst_pitch, src, 0, 0, width, height, kind, stream));
|
||||
}
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(size_t width, size_t height, hipMemcpyKind kind,
|
||||
bool enable_peer_access=false) {
|
||||
Memcpy2DFromArrayAsyncBenchmark benchmark;
|
||||
benchmark.AddSectionName("(" + std::to_string(width) + ", " + std::to_string(height) + ")");
|
||||
|
||||
const StreamGuard stream_guard(Streams::created);
|
||||
const hipStream_t stream = stream_guard.stream();
|
||||
|
||||
if (kind == hipMemcpyDeviceToHost) {
|
||||
size_t allocation_size = width * height * sizeof(int);
|
||||
LinearAllocGuard<int> host_allocation(LinearAllocs::hipHostMalloc, allocation_size);
|
||||
ArrayAllocGuard<int> array_allocation(make_hipExtent(width, height, 0), hipArrayDefault);
|
||||
benchmark.Run(host_allocation.ptr(), width * sizeof(int),
|
||||
array_allocation.ptr(), width * sizeof(int),
|
||||
height, hipMemcpyDeviceToHost, stream);
|
||||
} else {
|
||||
// hipMemcpyDeviceToDevice
|
||||
int src_device = std::get<0>(GetDeviceIds(enable_peer_access));
|
||||
int dst_device = std::get<1>(GetDeviceIds(enable_peer_access));
|
||||
|
||||
LinearAllocGuard2D<int> device_allocation(width, height);
|
||||
HIP_CHECK(hipSetDevice(dst_device));
|
||||
ArrayAllocGuard<int> array_allocation(make_hipExtent(width, height, 0), hipArrayDefault);
|
||||
HIP_CHECK(hipSetDevice(src_device));
|
||||
benchmark.Run(device_allocation.ptr(), device_allocation.pitch(),
|
||||
array_allocation.ptr(), device_allocation.width(),
|
||||
device_allocation.height(), hipMemcpyDeviceToDevice, stream);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy2DFromArrayAsync` from Device to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 8 KB x 32 B
|
||||
* - Large: 16 KB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy2DFromArrayAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy2DFromArrayAsync_DeviceToHost") {
|
||||
const auto width = GENERATE(4_KB, 8_KB, 16_KB);
|
||||
RunBenchmark(width, 32, hipMemcpyDeviceToHost);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy2DFromArrayAsync` from Device to Device with peer access disabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 8 KB x 32 B
|
||||
* - Large: 16 KB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy2DFromArrayAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy2DFromArrayAsync_DeviceToDevice_DisablePeerAccess") {
|
||||
const auto width = GENERATE(4_KB, 8_KB, 16_KB);
|
||||
RunBenchmark(width, 32, hipMemcpyDeviceToDevice);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy2DFromArrayAsync` from Device to Device with peer access enabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 8 KB x 32 B
|
||||
* - Large: 16 KB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy2DFromArrayAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Multi-device
|
||||
* - Device supports Peer-to-Peer access
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy2DFromArrayAsync_DeviceToDevice_EnablePeerAccess") {
|
||||
if (HipTest::getDeviceCount() < 2) {
|
||||
HipTest::HIP_SKIP_TEST("This test requires 2 GPUs. Skipping.");
|
||||
return;
|
||||
}
|
||||
const auto width = GENERATE(4_KB, 8_KB, 16_KB);
|
||||
RunBenchmark(width, 32, hipMemcpyDeviceToDevice, true);
|
||||
}
|
||||
@@ -0,0 +1,127 @@
|
||||
/*
|
||||
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 "memcpy_performance_common.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class Memcpy2DToArrayBenchmark : public Benchmark<Memcpy2DToArrayBenchmark> {
|
||||
public:
|
||||
void operator()(hipArray_t dst, const void* src, size_t src_pitch, size_t width,
|
||||
size_t height, hipMemcpyKind kind) {
|
||||
TIMED_SECTION(kTimerTypeCpu) {
|
||||
HIP_CHECK(hipMemcpy2DToArray(dst, 0, 0, src, src_pitch, width, height, kind));
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(size_t width, size_t height, hipMemcpyKind kind,
|
||||
bool enable_peer_access=false) {
|
||||
Memcpy2DToArrayBenchmark benchmark;
|
||||
benchmark.AddSectionName("(" + std::to_string(width) + ", " + std::to_string(height) + ")");
|
||||
|
||||
if (kind == hipMemcpyHostToDevice) {
|
||||
size_t allocation_size = width * height * sizeof(int);
|
||||
LinearAllocGuard<int> host_allocation(LinearAllocs::hipHostMalloc, allocation_size);
|
||||
ArrayAllocGuard<int> array_allocation(make_hipExtent(width, height, 0), hipArrayDefault);
|
||||
benchmark.Run(array_allocation.ptr(), host_allocation.ptr(), width * sizeof(int),
|
||||
width * sizeof(int), height, hipMemcpyHostToDevice);
|
||||
} else {
|
||||
// hipMemcpyDeviceToDevice
|
||||
int src_device = std::get<0>(GetDeviceIds(enable_peer_access));
|
||||
int dst_device = std::get<1>(GetDeviceIds(enable_peer_access));
|
||||
|
||||
LinearAllocGuard2D<int> device_allocation(width, height);
|
||||
HIP_CHECK(hipSetDevice(dst_device));
|
||||
ArrayAllocGuard<int> array_allocation(make_hipExtent(width, height, 0), hipArrayDefault);
|
||||
HIP_CHECK(hipSetDevice(src_device));
|
||||
benchmark.Run(array_allocation.ptr(), device_allocation.ptr(), device_allocation.pitch(),
|
||||
device_allocation.width(), device_allocation.height(), hipMemcpyDeviceToDevice);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy2DToArray` from Host to Device:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 8 KB x 32 B
|
||||
* - Large: 16 KB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy2DToArray.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy2DToArray_HostToDevice") {
|
||||
const auto width = GENERATE(4_KB, 8_KB, 16_KB);
|
||||
RunBenchmark(width, 32, hipMemcpyHostToDevice);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy2DToArray` from Device to Device with peer access disabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 8 KB x 32 B
|
||||
* - Large: 16 KB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy2DToArray.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy2DToArray_DeviceToDevice_DisablePeerAccess") {
|
||||
const auto width = GENERATE(4_KB, 8_KB, 16_KB);
|
||||
RunBenchmark(width, 32, hipMemcpyDeviceToDevice);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy2DToArray` from Device to Device with peer access enabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 8 KB x 32 B
|
||||
* - Large: 16 KB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy2DToArray.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Multi-device
|
||||
* - Device supports Peer-to-Peer access
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy2DToArray_DeviceToDevice_EnablePeerAccess") {
|
||||
if (HipTest::getDeviceCount() < 2) {
|
||||
HipTest::HIP_SKIP_TEST("This test requires 2 GPUs. Skipping.");
|
||||
return;
|
||||
}
|
||||
const auto width = GENERATE(4_KB, 8_KB, 16_KB);
|
||||
RunBenchmark(width, 32, hipMemcpyDeviceToDevice, true);
|
||||
}
|
||||
@@ -0,0 +1,133 @@
|
||||
/*
|
||||
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 "memcpy_performance_common.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class Memcpy2DToArrayAsyncBenchmark : public Benchmark<Memcpy2DToArrayAsyncBenchmark> {
|
||||
public:
|
||||
void operator()(hipArray_t dst, const void* src, size_t src_pitch, size_t width,
|
||||
size_t height, hipMemcpyKind kind, const hipStream_t& stream) {
|
||||
TIMED_SECTION_STREAM(kTimerTypeEvent, stream) {
|
||||
HIP_CHECK(hipMemcpy2DToArrayAsync(dst, 0, 0, src, src_pitch, width, height, kind, stream));
|
||||
}
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(size_t width, size_t height, hipMemcpyKind kind,
|
||||
bool enable_peer_access=false) {
|
||||
Memcpy2DToArrayAsyncBenchmark benchmark;
|
||||
benchmark.AddSectionName("(" + std::to_string(width) + ", " + std::to_string(height) + ")");
|
||||
|
||||
const StreamGuard stream_guard(Streams::created);
|
||||
const hipStream_t stream = stream_guard.stream();
|
||||
|
||||
if (kind == hipMemcpyHostToDevice) {
|
||||
size_t allocation_size = width * height * sizeof(int);
|
||||
LinearAllocGuard<int> host_allocation(LinearAllocs::hipHostMalloc, allocation_size);
|
||||
ArrayAllocGuard<int> array_allocation(make_hipExtent(width, height, 0), hipArrayDefault);
|
||||
benchmark.Run(array_allocation.ptr(), host_allocation.ptr(),
|
||||
width * sizeof(int), width * sizeof(int), height,
|
||||
hipMemcpyHostToDevice, stream);
|
||||
} else {
|
||||
// hipMemcpyDeviceToDevice
|
||||
int src_device = std::get<0>(GetDeviceIds(enable_peer_access));
|
||||
int dst_device = std::get<1>(GetDeviceIds(enable_peer_access));
|
||||
|
||||
LinearAllocGuard2D<int> device_allocation(width, height);
|
||||
HIP_CHECK(hipSetDevice(dst_device));
|
||||
ArrayAllocGuard<int> array_allocation(make_hipExtent(width, height, 0), hipArrayDefault);
|
||||
HIP_CHECK(hipSetDevice(src_device));
|
||||
benchmark.Run(array_allocation.ptr(), device_allocation.ptr(), device_allocation.pitch(),
|
||||
device_allocation.width(), device_allocation.height(),
|
||||
hipMemcpyDeviceToDevice, stream);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy2DToArrayAsync` from Host to Device:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 8 KB x 32 B
|
||||
* - Large: 16 KB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy2DToArrayAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy2DToArrayAsync_HostToDevice") {
|
||||
const auto width = GENERATE(4_KB, 8_KB, 16_KB);
|
||||
RunBenchmark(width, 32, hipMemcpyHostToDevice);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy2DToArrayAsync` from Device to Device with peer access disabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 8 KB x 32 B
|
||||
* - Large: 16 KB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy2DToArrayAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy2DToArrayAsync_DeviceToDevice_DisablePeerAccess") {
|
||||
const auto width = GENERATE(4_KB, 8_KB, 16_KB);
|
||||
RunBenchmark(width, 32, hipMemcpyDeviceToDevice);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy2DToArrayAsync` from Device to Device with peer access enabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 8 KB x 32 B
|
||||
* - Large: 16 KB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy2DToArrayAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Multi-device
|
||||
* - Device supports Peer-to-Peer access
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy2DToArrayAsync_DeviceToDevice_EnablePeerAccess") {
|
||||
if (HipTest::getDeviceCount() < 2) {
|
||||
HipTest::HIP_SKIP_TEST("This test requires 2 GPUs. Skipping.");
|
||||
return;
|
||||
}
|
||||
const auto width = GENERATE(4_KB, 8_KB, 16_KB);
|
||||
RunBenchmark(width, 32, hipMemcpyDeviceToDevice, true);
|
||||
}
|
||||
@@ -0,0 +1,189 @@
|
||||
/*
|
||||
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 "memcpy_performance_common.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class Memcpy3DBenchmark : public Benchmark<Memcpy3DBenchmark> {
|
||||
public:
|
||||
void operator()(const hipPitchedPtr& dst_ptr, const hipPitchedPtr& src_ptr,
|
||||
const hipExtent extent, hipMemcpyKind kind) {
|
||||
hipMemcpy3DParms params = CreateMemcpy3DParam(dst_ptr, make_hipPos(0, 0, 0),
|
||||
src_ptr, make_hipPos(0, 0, 0),
|
||||
extent, kind);
|
||||
TIMED_SECTION(kTimerTypeCpu) {
|
||||
HIP_CHECK(hipMemcpy3D(¶ms));
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(const hipExtent extent, hipMemcpyKind kind, bool enable_peer_access=false) {
|
||||
Memcpy3DBenchmark benchmark;
|
||||
benchmark.AddSectionName("(" + std::to_string(extent.width) + ", " + std::to_string(extent.height)
|
||||
+ ", " + std::to_string(extent.depth) + ")");
|
||||
|
||||
if (kind == hipMemcpyDeviceToHost) {
|
||||
LinearAllocGuard3D<int> device_allocation(extent);
|
||||
LinearAllocGuard<int> host_allocation(LinearAllocs::hipHostMalloc, device_allocation.width() *
|
||||
device_allocation.height() * device_allocation.depth());
|
||||
benchmark.Run(make_hipPitchedPtr(host_allocation.ptr(), device_allocation.width(),
|
||||
device_allocation.width(), device_allocation.height()),
|
||||
device_allocation.pitched_ptr(), device_allocation.extent(), kind);
|
||||
} else if (kind == hipMemcpyHostToDevice) {
|
||||
LinearAllocGuard3D<int> device_allocation(extent);
|
||||
LinearAllocGuard<int> host_allocation(LinearAllocs::hipHostMalloc, device_allocation.pitch() *
|
||||
device_allocation.height() * device_allocation.depth());
|
||||
benchmark.Run(device_allocation.pitched_ptr(),
|
||||
make_hipPitchedPtr(host_allocation.ptr(), device_allocation.pitch(),
|
||||
device_allocation.width(), device_allocation.height()),
|
||||
device_allocation.extent(), kind);
|
||||
} else if (kind == hipMemcpyHostToHost) {
|
||||
LinearAllocGuard3D<int> device_allocation(extent);
|
||||
LinearAllocGuard<int> src_allocation(LinearAllocs::hipHostMalloc, extent.width *
|
||||
extent.height * extent.depth);
|
||||
LinearAllocGuard<int> dst_allocation(LinearAllocs::hipHostMalloc, extent.width *
|
||||
extent.height * extent.depth);
|
||||
benchmark.Run(make_hipPitchedPtr(dst_allocation.ptr(), extent.width, extent.width, extent.height),
|
||||
make_hipPitchedPtr(src_allocation.ptr(), extent.width, extent.width, extent.height),
|
||||
extent, kind);
|
||||
} else {
|
||||
// hipMemcpyDeviceToDevice
|
||||
int src_device = std::get<0>(GetDeviceIds(enable_peer_access));
|
||||
int dst_device = std::get<1>(GetDeviceIds(enable_peer_access));
|
||||
|
||||
LinearAllocGuard3D<int> src_allocation(extent);
|
||||
HIP_CHECK(hipSetDevice(dst_device));
|
||||
LinearAllocGuard3D<int> dst_allocation(extent);
|
||||
|
||||
HIP_CHECK(hipSetDevice(src_device));
|
||||
benchmark.Run(dst_allocation.pitched_ptr(), src_allocation.pitched_ptr(),
|
||||
dst_allocation.extent(), kind);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy3D` from Device to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 16 B x 4 B
|
||||
* - Medium: 4 MB x 16 B x 4 B
|
||||
* - Large: 16 MB x 16 B x 4 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy3D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy3D_DeviceToHost") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(make_hipExtent(width, 16, 4), hipMemcpyDeviceToHost);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy3D` from Host to Device:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 16 B x 4 B
|
||||
* - Medium: 4 MB x 16 B x 4 B
|
||||
* - Large: 16 MB x 16 B x 4 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy3D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy3D_HostToDevice") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(make_hipExtent(width, 16, 4), hipMemcpyHostToDevice);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy3D` from Host to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 16 B x 4 B
|
||||
* - Medium: 4 MB x 16 B x 4 B
|
||||
* - Large: 16 MB x 16 B x 4 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy3D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy3D_HostToHost") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(make_hipExtent(width, 16, 4), hipMemcpyHostToHost);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy3D` from Device to Device with peer access disabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 16 B x 4 B
|
||||
* - Medium: 4 MB x 16 B x 4 B
|
||||
* - Large: 16 MB x 16 B x 4 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy3D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy3D_DeviceToDevice_DisablePeerAccess") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(make_hipExtent(width, 16, 4), hipMemcpyDeviceToDevice);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy3D` from Device to Device with peer access enabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 16 B x 4 B
|
||||
* - Medium: 4 MB x 16 B x 4 B
|
||||
* - Large: 16 MB x 16 B x 4 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy3D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Multi-device
|
||||
* - Device supports Peer-to-Peer access
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy3D_DeviceToDevice_EnablePeerAccess") {
|
||||
if (HipTest::getDeviceCount() < 2) {
|
||||
HipTest::HIP_SKIP_TEST("This test requires 2 GPUs. Skipping.");
|
||||
return;
|
||||
}
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(make_hipExtent(width, 16, 4), hipMemcpyDeviceToDevice, true);
|
||||
}
|
||||
@@ -0,0 +1,192 @@
|
||||
/*
|
||||
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 "memcpy_performance_common.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class Memcpy3DAsyncBenchmark : public Benchmark<Memcpy3DAsyncBenchmark> {
|
||||
public:
|
||||
void operator()(const hipPitchedPtr& dst_ptr, const hipPitchedPtr& src_ptr,
|
||||
const hipExtent extent, hipMemcpyKind kind, const hipStream_t& stream) {
|
||||
hipMemcpy3DParms params = CreateMemcpy3DParam(dst_ptr, make_hipPos(0, 0, 0),
|
||||
src_ptr, make_hipPos(0, 0, 0),
|
||||
extent, kind);
|
||||
TIMED_SECTION_STREAM(kTimerTypeEvent, stream) {
|
||||
HIP_CHECK(hipMemcpy3DAsync(¶ms, stream));
|
||||
}
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(const hipExtent extent, hipMemcpyKind kind, bool enable_peer_access=false) {
|
||||
Memcpy3DAsyncBenchmark benchmark;
|
||||
benchmark.AddSectionName("(" + std::to_string(extent.width) + ", " + std::to_string(extent.height)
|
||||
+ ", " + std::to_string(extent.depth) + ")");
|
||||
|
||||
const StreamGuard stream_guard(Streams::created);
|
||||
const hipStream_t stream = stream_guard.stream();
|
||||
|
||||
if (kind == hipMemcpyDeviceToHost) {
|
||||
LinearAllocGuard3D<int> device_allocation(extent);
|
||||
LinearAllocGuard<int> host_allocation(LinearAllocs::hipHostMalloc, device_allocation.width() *
|
||||
device_allocation.height() * device_allocation.depth());
|
||||
benchmark.Run(make_hipPitchedPtr(host_allocation.ptr(), device_allocation.width(),
|
||||
device_allocation.width(), device_allocation.height()),
|
||||
device_allocation.pitched_ptr(), device_allocation.extent(), kind, stream);
|
||||
} else if (kind == hipMemcpyHostToDevice) {
|
||||
LinearAllocGuard3D<int> device_allocation(extent);
|
||||
LinearAllocGuard<int> host_allocation(LinearAllocs::hipHostMalloc, device_allocation.pitch() *
|
||||
device_allocation.height() * device_allocation.depth());
|
||||
benchmark.Run(device_allocation.pitched_ptr(),
|
||||
make_hipPitchedPtr(host_allocation.ptr(),
|
||||
device_allocation.pitch(),
|
||||
device_allocation.width(),
|
||||
device_allocation.height()),
|
||||
device_allocation.extent(), kind, stream);
|
||||
} else if (kind == hipMemcpyHostToHost) {
|
||||
LinearAllocGuard3D<int> device_allocation(extent);
|
||||
LinearAllocGuard<int> src_allocation(LinearAllocs::hipHostMalloc, extent.width *
|
||||
extent.height * extent.depth);
|
||||
LinearAllocGuard<int> dst_allocation(LinearAllocs::hipHostMalloc, extent.width *
|
||||
extent.height * extent.depth);
|
||||
benchmark.Run(make_hipPitchedPtr(dst_allocation.ptr(), extent.width, extent.width, extent.height),
|
||||
make_hipPitchedPtr(src_allocation.ptr(), extent.width, extent.width, extent.height),
|
||||
extent, kind, stream);
|
||||
} else {
|
||||
// hipMemcpyDeviceToDevice
|
||||
int src_device = std::get<0>(GetDeviceIds(enable_peer_access));
|
||||
int dst_device = std::get<1>(GetDeviceIds(enable_peer_access));
|
||||
|
||||
LinearAllocGuard3D<int> src_allocation(extent);
|
||||
HIP_CHECK(hipSetDevice(dst_device));
|
||||
LinearAllocGuard3D<int> dst_allocation(extent);
|
||||
HIP_CHECK(hipSetDevice(src_device));
|
||||
benchmark.Run(dst_allocation.pitched_ptr(), src_allocation.pitched_ptr(),
|
||||
dst_allocation.extent(), kind, stream);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy3DAsync` from Device to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 16 B x 4 B
|
||||
* - Medium: 4 MB x 16 B x 4 B
|
||||
* - Large: 16 MB x 16 B x 4 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy3DAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy3DAsync_DeviceToHost") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(make_hipExtent(width, 16, 4), hipMemcpyDeviceToHost);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy3DAsync` from Host to Device:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 16 B x 4 B
|
||||
* - Medium: 4 MB x 16 B x 4 B
|
||||
* - Large: 16 MB x 16 B x 4 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy3DAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy3DAsync_HostToDevice") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(make_hipExtent(width, 16, 4), hipMemcpyHostToDevice);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy3DAsync` from Host to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 16 B x 4 B
|
||||
* - Medium: 4 MB x 16 B x 4 B
|
||||
* - Large: 16 MB x 16 B x 4 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy3DAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy3DAsync_HostToHost") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(make_hipExtent(width, 16, 4), hipMemcpyHostToHost);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy3DAsync` from Device to Device with peer access disabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 16 B x 4 B
|
||||
* - Medium: 4 MB x 16 B x 4 B
|
||||
* - Large: 16 MB x 16 B x 4 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy3DAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy3DAsync_DeviceToDevice_DisablePeerAccess") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(make_hipExtent(width, 16, 4), hipMemcpyDeviceToDevice);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy3DAsync` from Device to Device with peer access enabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 16 B x 4 B
|
||||
* - Medium: 4 MB x 16 B x 4 B
|
||||
* - Large: 16 MB x 16 B x 4 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpy3DAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpy3DAsync_DeviceToDevice_EnablePeerAccess") {
|
||||
if (HipTest::getDeviceCount() < 2) {
|
||||
HipTest::HIP_SKIP_TEST("This test requires 2 GPUs. Skipping.");
|
||||
return;
|
||||
}
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(make_hipExtent(width, 16, 4), hipMemcpyDeviceToDevice, true);
|
||||
}
|
||||
@@ -0,0 +1,192 @@
|
||||
/*
|
||||
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 "memcpy_performance_common.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class MemcpyAsyncBenchmark : public Benchmark<MemcpyAsyncBenchmark> {
|
||||
public:
|
||||
void operator()(void* dst, const void* src, size_t size, hipMemcpyKind kind, const hipStream_t& stream) {
|
||||
TIMED_SECTION_STREAM(kTimerTypeEvent, stream) {
|
||||
HIP_CHECK(hipMemcpyAsync(dst, src, size, kind, stream));
|
||||
}
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(LinearAllocs dst_allocation_type, LinearAllocs src_allocation_type,
|
||||
size_t size, hipMemcpyKind kind, bool enable_peer_access=false) {
|
||||
MemcpyAsyncBenchmark benchmark;
|
||||
benchmark.AddSectionName(std::to_string(size));
|
||||
benchmark.AddSectionName(GetAllocationSectionName(src_allocation_type));
|
||||
benchmark.AddSectionName(GetAllocationSectionName(dst_allocation_type));
|
||||
|
||||
const StreamGuard stream_guard{Streams::created};
|
||||
const hipStream_t stream = stream_guard.stream();
|
||||
if (kind != hipMemcpyDeviceToDevice) {
|
||||
LinearAllocGuard<int> src_allocation(src_allocation_type, size);
|
||||
LinearAllocGuard<int> dst_allocation(dst_allocation_type, size);
|
||||
benchmark.Run(dst_allocation.ptr(), src_allocation.ptr(), size, kind, stream);
|
||||
} else {
|
||||
int src_device = std::get<0>(GetDeviceIds(enable_peer_access));
|
||||
int dst_device = std::get<1>(GetDeviceIds(enable_peer_access));
|
||||
|
||||
LinearAllocGuard<int> src_allocation(src_allocation_type, size);
|
||||
HIP_CHECK(hipSetDevice(dst_device));
|
||||
LinearAllocGuard<int> dst_allocation(dst_allocation_type, size);
|
||||
HIP_CHECK(hipSetDevice(src_device));
|
||||
benchmark.Run(dst_allocation.ptr(), src_allocation.ptr(), size, kind, stream);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyAsync` from Device to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: device malloc
|
||||
* - Destination: host pinned and pageable
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyAsync_DeviceToHost") {
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto src_allocation_type = LinearAllocs::hipMalloc;
|
||||
const auto dst_allocation_type = GENERATE(LinearAllocs::malloc, LinearAllocs::hipHostMalloc);
|
||||
RunBenchmark(dst_allocation_type, src_allocation_type, allocation_size, hipMemcpyDeviceToHost);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyAsync` from Host to Device:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: host pinned and pageable
|
||||
* - Destination: device malloc
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyAsync_HostToDevice") {
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto src_allocation_type = GENERATE(LinearAllocs::malloc, LinearAllocs::hipHostMalloc);
|
||||
const auto dst_allocation_type = LinearAllocs::hipMalloc;
|
||||
RunBenchmark(dst_allocation_type, src_allocation_type, allocation_size, hipMemcpyHostToDevice);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyAsync` from Host to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: host pinned and pageable
|
||||
* - Destination: host pinned and pageable
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyAsync_HostToHost") {
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto src_allocation_type = GENERATE(LinearAllocs::malloc, LinearAllocs::hipHostMalloc);
|
||||
const auto dst_allocation_type = GENERATE(LinearAllocs::malloc, LinearAllocs::hipHostMalloc);
|
||||
RunBenchmark(dst_allocation_type, src_allocation_type, allocation_size, hipMemcpyHostToHost);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyAsync` from Device to Device with peer access disabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: device malloc
|
||||
* - Destination: device malloc
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyAsync_DeviceToDevice_DisablePeerAccess") {
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto src_allocation_type = LinearAllocs::hipMalloc;
|
||||
const auto dst_allocation_type = LinearAllocs::hipMalloc;
|
||||
RunBenchmark(dst_allocation_type, src_allocation_type, allocation_size, hipMemcpyDeviceToDevice);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyAsync` from Device to Device with peer access enabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: device malloc
|
||||
* - Destination: device malloc
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Multi-device
|
||||
* - Device supports Peer-to-Peer access
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyAsync_DeviceToDevice_EnablePeerAccess") {
|
||||
if (HipTest::getDeviceCount() < 2) {
|
||||
HipTest::HIP_SKIP_TEST("This test requires 2 GPUs. Skipping.");
|
||||
return;
|
||||
}
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto src_allocation_type = LinearAllocs::hipMalloc;
|
||||
const auto dst_allocation_type = LinearAllocs::hipMalloc;
|
||||
RunBenchmark(dst_allocation_type, src_allocation_type, allocation_size, hipMemcpyDeviceToDevice, true);
|
||||
}
|
||||
@@ -0,0 +1,69 @@
|
||||
/*
|
||||
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 "memcpy_performance_common.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class MemcpyAtoHBenchmark : public Benchmark<MemcpyAtoHBenchmark> {
|
||||
public:
|
||||
void operator()(void* dst, hipArray_t src_array, size_t allocation_size) {
|
||||
TIMED_SECTION(kTimerTypeCpu) {
|
||||
HIP_CHECK(hipMemcpyAtoH(dst, src_array, 0, allocation_size));
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(LinearAllocs host_allocation_type, size_t width) {
|
||||
MemcpyAtoHBenchmark benchmark;
|
||||
benchmark.AddSectionName(std::to_string(width));
|
||||
benchmark.AddSectionName(GetAllocationSectionName(host_allocation_type));
|
||||
|
||||
size_t allocation_size = width * sizeof(int);
|
||||
LinearAllocGuard<int> host_allocation(host_allocation_type, allocation_size);
|
||||
ArrayAllocGuard<int> array_allocation(make_hipExtent(width, 0, 0), hipArrayDefault);
|
||||
benchmark.Run(host_allocation.ptr(), array_allocation.ptr(), allocation_size);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyAtoH` from Device array to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 512 B
|
||||
* - Medium: 1024 B
|
||||
* - Large: 4096 B
|
||||
* -# Allocation type
|
||||
* - Host: host pinned and pageable
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyAtoH.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyAtoH") {
|
||||
const auto allocation_size = GENERATE(512, 1024, 4096);
|
||||
const auto host_allocation_type = GENERATE(LinearAllocs::malloc, LinearAllocs::hipHostMalloc);
|
||||
RunBenchmark(host_allocation_type, allocation_size);
|
||||
}
|
||||
@@ -0,0 +1,103 @@
|
||||
/*
|
||||
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 "memcpy_performance_common.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class MemcpyDtoDBenchmark : public Benchmark<MemcpyDtoDBenchmark> {
|
||||
public:
|
||||
void operator()(hipDeviceptr_t& dst, const hipDeviceptr_t& src, size_t size) {
|
||||
TIMED_SECTION(kTimerTypeCpu) {
|
||||
HIP_CHECK(hipMemcpyDtoD(dst, src, size));
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(size_t size, bool enable_peer_access=false) {
|
||||
MemcpyDtoDBenchmark benchmark;
|
||||
benchmark.AddSectionName(std::to_string(size));
|
||||
|
||||
int src_device = std::get<0>(GetDeviceIds(enable_peer_access));
|
||||
int dst_device = std::get<1>(GetDeviceIds(enable_peer_access));
|
||||
|
||||
LinearAllocGuard<int> src_allocation(LinearAllocs::hipMalloc, size);
|
||||
HIP_CHECK(hipSetDevice(dst_device));
|
||||
LinearAllocGuard<int> dst_allocation(LinearAllocs::hipMalloc, size);
|
||||
HIP_CHECK(hipSetDevice(src_device));
|
||||
|
||||
benchmark.Run(reinterpret_cast<hipDeviceptr_t>(dst_allocation.ptr()),
|
||||
reinterpret_cast<hipDeviceptr_t>(src_allocation.ptr()), size);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyDtoD` from Device to Device with peer access enabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: device malloc
|
||||
* - Destination: device malloc
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyDtoD.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Multi-device
|
||||
* - Device supports Peer-to-Peer access
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyDtoD_PeerAccessEnabled") {
|
||||
if (HipTest::getDeviceCount() < 2) {
|
||||
HipTest::HIP_SKIP_TEST("This test requires 2 GPUs. Skipping.");
|
||||
return;
|
||||
}
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(allocation_size, true);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyDtoD` from Device to Device with peer access disabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: device malloc
|
||||
* - Destination: device malloc
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyDtoD.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyDtoD_PeerAccessDisabled") {
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(allocation_size);
|
||||
}
|
||||
@@ -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 "memcpy_performance_common.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class MemcpyDtoDAsyncBenchmark : public Benchmark<MemcpyDtoDAsyncBenchmark> {
|
||||
public:
|
||||
void operator()(hipDeviceptr_t& dst, const hipDeviceptr_t& src, size_t size, const hipStream_t& stream) {
|
||||
TIMED_SECTION_STREAM(kTimerTypeEvent, stream) {
|
||||
HIP_CHECK(hipMemcpyDtoDAsync(dst, src, size, stream));
|
||||
}
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(size_t size, bool enable_peer_access=false) {
|
||||
MemcpyDtoDAsyncBenchmark benchmark;
|
||||
benchmark.AddSectionName(std::to_string(size));
|
||||
|
||||
const StreamGuard stream_guard(Streams::created);
|
||||
const hipStream_t stream = stream_guard.stream();
|
||||
int src_device = std::get<0>(GetDeviceIds(enable_peer_access));
|
||||
int dst_device = std::get<1>(GetDeviceIds(enable_peer_access));
|
||||
|
||||
LinearAllocGuard<int> src_allocation(LinearAllocs::hipMalloc, size);
|
||||
HIP_CHECK(hipSetDevice(dst_device));
|
||||
LinearAllocGuard<int> dst_allocation(LinearAllocs::hipMalloc, size);
|
||||
HIP_CHECK(hipSetDevice(src_device));
|
||||
benchmark.Run(reinterpret_cast<hipDeviceptr_t>(dst_allocation.ptr()),
|
||||
reinterpret_cast<hipDeviceptr_t>(src_allocation.ptr()),
|
||||
size, stream);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyDtoDAsync` from Device to Device with peer access enabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: device malloc
|
||||
* - Destination: device malloc
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyDtoDAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Multi-device
|
||||
* - Device supports Peer-to-Peer access
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyDtoDAsync_PeerAccessEnabled") {
|
||||
if (HipTest::getDeviceCount() < 2) {
|
||||
HipTest::HIP_SKIP_TEST("This test requires 2 GPUs. Skipping.");
|
||||
return;
|
||||
}
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(allocation_size, true);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyDtoD` from Device to Device with peer access disabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: device malloc
|
||||
* - Destination: device malloc
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyDtoDAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyDtoDAsync_PeerAccessDisabled") {
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(allocation_size);
|
||||
}
|
||||
@@ -0,0 +1,72 @@
|
||||
/*
|
||||
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 "memcpy_performance_common.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class MemcpyDtoHBenchmark : public Benchmark<MemcpyDtoHBenchmark> {
|
||||
public:
|
||||
void operator()(void* dst, const hipDeviceptr_t& src, size_t size) {
|
||||
TIMED_SECTION(kTimerTypeCpu) {
|
||||
HIP_CHECK(hipMemcpyDtoH(dst, src, size));
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(LinearAllocs host_allocation_type, LinearAllocs device_allocation_type, size_t size) {
|
||||
MemcpyDtoHBenchmark benchmark;
|
||||
benchmark.AddSectionName(std::to_string(size));
|
||||
benchmark.AddSectionName(GetAllocationSectionName(host_allocation_type));
|
||||
|
||||
LinearAllocGuard<int> device_allocation(device_allocation_type, size);
|
||||
LinearAllocGuard<int> host_allocation(host_allocation_type, size);
|
||||
benchmark.Run(host_allocation.ptr(),
|
||||
reinterpret_cast<hipDeviceptr_t>(device_allocation.ptr()),
|
||||
size);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyDtoH` from Device to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: device malloc
|
||||
* - Destination: host pinned and pageable
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyDtoH.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyDtoH") {
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto device_allocation_type = LinearAllocs::hipMalloc;
|
||||
const auto host_allocation_type = GENERATE(LinearAllocs::malloc, LinearAllocs::hipHostMalloc);
|
||||
RunBenchmark(host_allocation_type, device_allocation_type, allocation_size);
|
||||
}
|
||||
@@ -0,0 +1,75 @@
|
||||
/*
|
||||
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 "memcpy_performance_common.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class MemcpyDtoHAsyncBenchmark : public Benchmark<MemcpyDtoHAsyncBenchmark> {
|
||||
public:
|
||||
void operator()(void* dst, const hipDeviceptr_t& src, size_t size, const hipStream_t& stream) {
|
||||
TIMED_SECTION_STREAM(kTimerTypeEvent, stream) {
|
||||
HIP_CHECK(hipMemcpyDtoHAsync(dst, src, size, stream));
|
||||
}
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(LinearAllocs host_allocation_type, LinearAllocs device_allocation_type, size_t size) {
|
||||
MemcpyDtoHAsyncBenchmark benchmark;
|
||||
benchmark.AddSectionName(std::to_string(size));
|
||||
benchmark.AddSectionName(GetAllocationSectionName(host_allocation_type));
|
||||
|
||||
const StreamGuard stream_guard(Streams::created);
|
||||
const hipStream_t stream = stream_guard.stream();
|
||||
LinearAllocGuard<int> device_allocation(device_allocation_type, size);
|
||||
LinearAllocGuard<int> host_allocation(host_allocation_type, size);
|
||||
benchmark.Run(host_allocation.ptr(),
|
||||
reinterpret_cast<hipDeviceptr_t>(device_allocation.ptr()),
|
||||
size, stream);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyDtoHAsync` from Device to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: device malloc
|
||||
* - Destination: host pinned and pageable
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyDtoHAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyDtoHAsync") {
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto device_allocation_type = LinearAllocs::hipMalloc;
|
||||
const auto host_allocation_type = GENERATE(LinearAllocs::malloc, LinearAllocs::hipHostMalloc);
|
||||
RunBenchmark(host_allocation_type, device_allocation_type, allocation_size);
|
||||
}
|
||||
@@ -0,0 +1,116 @@
|
||||
/*
|
||||
Copyright (c) 2024 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 "memcpy_performance_common.hh"
|
||||
#pragma clang diagnostic ignored "-Wvla-extension"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
__device__ int devSymbol[1_MB];
|
||||
|
||||
class MemcpyFromSymbolBenchmark : public Benchmark<MemcpyFromSymbolBenchmark> {
|
||||
public:
|
||||
void operator()(const void* source, void* result, size_t size, size_t offset) {
|
||||
HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), source, size, offset));
|
||||
TIMED_SECTION(kTimerTypeCpu) {
|
||||
HIP_CHECK(hipMemcpyFromSymbol(result, HIP_SYMBOL(devSymbol), size, offset));
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(const void* source, void* result, size_t size=1, size_t offset=0) {
|
||||
MemcpyFromSymbolBenchmark benchmark;
|
||||
benchmark.AddSectionName(std::to_string(size));
|
||||
benchmark.AddSectionName(std::to_string(offset));
|
||||
benchmark.Run(source, result, size, offset);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyFromSymbol` from Device to Host.
|
||||
* - Utilizes sigular integer values.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyFromSymbol.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyFromSymbol_SingularValue") {
|
||||
int set{42};
|
||||
int result{0};
|
||||
RunBenchmark(&set, &result);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyFromSymbol` from Device to Host.
|
||||
* - Utilizes array integers:
|
||||
* - Small: 1 KB
|
||||
* - Medium: 4 KB
|
||||
* - Large: 512 KB
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyFromSymbol.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyFromSymbol_ArrayValue") {
|
||||
size_t size = GENERATE(1_KB, 4_KB, 512_KB);
|
||||
std::vector<int> array(size);
|
||||
std::fill_n(array.data(), size, 42);
|
||||
std::vector<int> result(size);
|
||||
std::fill_n(result.data(), size, 0);
|
||||
|
||||
RunBenchmark(array.data(), result.data(), sizeof(int) * size);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyFromSymbol` from Device to Host.
|
||||
* - Utilizes array integers with offsets:
|
||||
* - Small: 1 KB
|
||||
* - Medium: 4 KB
|
||||
* - Large: 512 KB
|
||||
* - Offset: 0 and size/2
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyFromSymbol.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyFromSymbol_WithOffset") {
|
||||
size_t size = GENERATE(1_KB, 4_KB, 512_KB);
|
||||
std::vector<int> array(size);
|
||||
std::fill_n(array.data(), size, 42);
|
||||
std::vector<int> result(size);
|
||||
std::fill_n(result.data(), size, 0);
|
||||
|
||||
size_t offset = GENERATE_REF(0, size / 2);
|
||||
RunBenchmark(array.data() + offset, result.data() + offset, sizeof(int) * (size - offset), offset * sizeof(int));
|
||||
}
|
||||
@@ -0,0 +1,122 @@
|
||||
/*
|
||||
Copyright (c) 2024 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 "memcpy_performance_common.hh"
|
||||
#pragma clang diagnostic ignored "-Wvla-extension"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
__device__ int devSymbol[1_MB];
|
||||
|
||||
class MemcpyFromSymbolAsyncBenchmark : public Benchmark<MemcpyFromSymbolAsyncBenchmark> {
|
||||
public:
|
||||
void operator()(const void* source, void* result, size_t size, size_t offset, const hipStream_t& stream) {
|
||||
HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), source, size, offset,
|
||||
hipMemcpyHostToDevice, stream));
|
||||
TIMED_SECTION_STREAM(kTimerTypeEvent, stream) {
|
||||
HIP_CHECK(hipMemcpyFromSymbolAsync(result, HIP_SYMBOL(devSymbol), size, offset,
|
||||
hipMemcpyDeviceToHost, stream));
|
||||
}
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(const void* source, void* result, size_t size=1, size_t offset=0) {
|
||||
MemcpyFromSymbolAsyncBenchmark benchmark;
|
||||
benchmark.AddSectionName(std::to_string(size));
|
||||
benchmark.AddSectionName(std::to_string(offset));
|
||||
|
||||
const StreamGuard stream_guard(Streams::created);
|
||||
const hipStream_t stream = stream_guard.stream();
|
||||
benchmark.Run(source, result, size, offset, stream);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyFromSymbolAsync` from Device to Host.
|
||||
* - Utilizes sigular integer values.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyFromSymbolAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyFromSymbolAsync_SingularValue") {
|
||||
int set{42};
|
||||
int result{0};
|
||||
RunBenchmark(&set, &result);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyFromSymbolAsync` from Device to Host.
|
||||
* - Utilizes array integers:
|
||||
* - Small: 1 KB
|
||||
* - Medium: 4 KB
|
||||
* - Large: 512 KB
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyFromSymbolAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyFromSymbolAsync_ArrayValue") {
|
||||
size_t size = GENERATE(1_KB, 4_KB, 512_KB);
|
||||
std::vector<int> array(size);
|
||||
std::fill_n(array.data(), size, 42);
|
||||
std::vector<int> result(size);
|
||||
std::fill_n(result.data(), size, 0);
|
||||
|
||||
RunBenchmark(array.data(), result.data(), sizeof(int) * size);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyFromSymbolAsync` from Device to Host.
|
||||
* - Utilizes array integers with offsets:
|
||||
* - Small: 1 KB
|
||||
* - Medium: 4 KB
|
||||
* - Large: 512 KB
|
||||
* - Offset: 0 and size/2
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyFromSymbolAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyFromSymbolAsync_WithOffset") {
|
||||
size_t size = GENERATE(1_KB, 4_KB, 512_KB);
|
||||
std::vector<int> array(size);
|
||||
std::fill_n(array.data(), size, 42);
|
||||
std::vector<int> result(size);
|
||||
std::fill_n(result.data(), size, 0);
|
||||
|
||||
size_t offset = GENERATE_REF(0, size / 2);
|
||||
RunBenchmark(array.data() + offset, result.data() + offset, sizeof(int) * (size - offset), offset * sizeof(int));
|
||||
}
|
||||
@@ -0,0 +1,69 @@
|
||||
/*
|
||||
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 "memcpy_performance_common.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class MemcpyHtoABenchmark : public Benchmark<MemcpyHtoABenchmark> {
|
||||
public:
|
||||
void operator()(hipArray_t dst_array, const void* src, size_t allocation_size) {
|
||||
TIMED_SECTION(kTimerTypeCpu) {
|
||||
HIP_CHECK(hipMemcpyHtoA(dst_array, 0, src, allocation_size));
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(LinearAllocs host_allocation_type, size_t width) {
|
||||
MemcpyHtoABenchmark benchmark;
|
||||
benchmark.AddSectionName(std::to_string(width));
|
||||
benchmark.AddSectionName(GetAllocationSectionName(host_allocation_type));
|
||||
|
||||
size_t allocation_size = width * sizeof(int);
|
||||
ArrayAllocGuard<int> array_allocation(make_hipExtent(width, 0, 0), hipArrayDefault);
|
||||
LinearAllocGuard<int> host_allocation(host_allocation_type, allocation_size);
|
||||
benchmark.Run(array_allocation.ptr(), host_allocation.ptr(), allocation_size);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyHtoA` from Host to Device array:
|
||||
* -# Allocation size
|
||||
* - Small: 512 B
|
||||
* - Medium: 1024 B
|
||||
* - Large: 4096 B
|
||||
* -# Allocation type
|
||||
* - Host: host pinned and pageable
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyHtoA.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyHtoA") {
|
||||
const auto allocation_size = GENERATE(512, 1024, 4096);
|
||||
const auto host_allocation_type = GENERATE(LinearAllocs::malloc, LinearAllocs::hipHostMalloc);
|
||||
RunBenchmark(host_allocation_type, allocation_size);
|
||||
}
|
||||
@@ -0,0 +1,70 @@
|
||||
/*
|
||||
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 "memcpy_performance_common.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class MemcpyHtoDBenchmark : public Benchmark<MemcpyHtoDBenchmark> {
|
||||
public:
|
||||
void operator()(hipDeviceptr_t& dst, void* src, size_t size) {
|
||||
TIMED_SECTION(kTimerTypeCpu) {
|
||||
HIP_CHECK(hipMemcpyHtoD(dst, src, size));
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(LinearAllocs host_allocation_type, LinearAllocs device_allocation_type, size_t size) {
|
||||
MemcpyHtoDBenchmark benchmark;
|
||||
benchmark.AddSectionName(std::to_string(size));
|
||||
benchmark.AddSectionName(GetAllocationSectionName(host_allocation_type));
|
||||
|
||||
LinearAllocGuard<int> device_allocation(device_allocation_type, size);
|
||||
LinearAllocGuard<int> host_allocation(host_allocation_type, size);
|
||||
benchmark.Run(reinterpret_cast<hipDeviceptr_t>(device_allocation.ptr()), host_allocation.ptr(), size);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyHtoD` from Host to Device:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: host pinned and pageable
|
||||
* - Destination: device malloc
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyHtoD.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyHtoD") {
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto device_allocation_type = LinearAllocs::hipMalloc;
|
||||
const auto host_allocation_type = GENERATE(LinearAllocs::malloc, LinearAllocs::hipHostMalloc);
|
||||
RunBenchmark(host_allocation_type, device_allocation_type, allocation_size);
|
||||
}
|
||||
@@ -0,0 +1,74 @@
|
||||
/*
|
||||
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 "memcpy_performance_common.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class MemcpyHtoDAsyncBenchmark : public Benchmark<MemcpyHtoDAsyncBenchmark> {
|
||||
public:
|
||||
void operator()(hipDeviceptr_t& dst, void* src, size_t size, const hipStream_t& stream) {
|
||||
TIMED_SECTION_STREAM(kTimerTypeEvent, stream) {
|
||||
HIP_CHECK(hipMemcpyHtoDAsync(dst, src, size, stream));
|
||||
}
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(LinearAllocs host_allocation_type, LinearAllocs device_allocation_type, size_t size) {
|
||||
MemcpyHtoDAsyncBenchmark benchmark;
|
||||
benchmark.AddSectionName(std::to_string(size));
|
||||
benchmark.AddSectionName(GetAllocationSectionName(host_allocation_type));
|
||||
|
||||
const StreamGuard stream_guard(Streams::created);
|
||||
const hipStream_t stream = stream_guard.stream();
|
||||
LinearAllocGuard<int> device_allocation(device_allocation_type, size);
|
||||
LinearAllocGuard<int> host_allocation(host_allocation_type, size);
|
||||
benchmark.Run(reinterpret_cast<hipDeviceptr_t>(device_allocation.ptr()),
|
||||
host_allocation.ptr(), size, stream);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyHtoD` from Host to Device:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: host pinned and pageable
|
||||
* - Destination: device malloc
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyHtoDAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyHtoDAsync") {
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto device_allocation_type = LinearAllocs::hipMalloc;
|
||||
const auto host_allocation_type = GENERATE(LinearAllocs::malloc, LinearAllocs::hipHostMalloc);
|
||||
RunBenchmark(host_allocation_type, device_allocation_type, allocation_size);
|
||||
}
|
||||
@@ -0,0 +1,188 @@
|
||||
/*
|
||||
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 "memcpy_performance_common.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class MemcpyParam2DBenchmark : public Benchmark<MemcpyParam2DBenchmark> {
|
||||
public:
|
||||
void operator()(void* dst, size_t dst_pitch, void* src, size_t src_pitch,
|
||||
size_t width, size_t height, hipMemcpyKind kind) {
|
||||
hip_Memcpy2D params = CreateMemcpy2DParam(dst, dst_pitch, src, src_pitch,
|
||||
width, height, kind);
|
||||
TIMED_SECTION(kTimerTypeCpu) {
|
||||
HIP_CHECK(hipMemcpyParam2D(¶ms));
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(size_t width, size_t height, hipMemcpyKind kind,
|
||||
bool enable_peer_access=false) {
|
||||
MemcpyParam2DBenchmark benchmark;
|
||||
benchmark.AddSectionName("(" + std::to_string(width) + ", " + std::to_string(height) + ")");
|
||||
|
||||
if (kind == hipMemcpyDeviceToHost) {
|
||||
LinearAllocGuard2D<int> device_allocation(width, height);
|
||||
LinearAllocGuard<int> host_allocation(LinearAllocs::hipHostMalloc, device_allocation.width() * height);
|
||||
benchmark.Run(host_allocation.ptr(), device_allocation.width(),
|
||||
device_allocation.ptr(), device_allocation.pitch(),
|
||||
device_allocation.width(), device_allocation.height(), kind);
|
||||
} else if (kind == hipMemcpyHostToDevice) {
|
||||
LinearAllocGuard2D<int> device_allocation(width, height);
|
||||
LinearAllocGuard<int> host_allocation(LinearAllocs::hipHostMalloc, device_allocation.width() * height);
|
||||
benchmark.Run(device_allocation.ptr(), device_allocation.pitch(),
|
||||
host_allocation.ptr(), device_allocation.width(),
|
||||
device_allocation.width(), device_allocation.height(), kind);
|
||||
} else if (kind == hipMemcpyHostToHost) {
|
||||
LinearAllocGuard<int> src_allocation(LinearAllocs::hipHostMalloc, width * sizeof(int) * height);
|
||||
LinearAllocGuard<int> dst_allocation(LinearAllocs::hipHostMalloc, width * sizeof(int) * height);
|
||||
benchmark.Run(dst_allocation.ptr(), width * sizeof(int),
|
||||
src_allocation.ptr(), width * sizeof(int),
|
||||
width * sizeof(int), height, kind);
|
||||
} else {
|
||||
// hipMemcpyDeviceToDevice
|
||||
int src_device = std::get<0>(GetDeviceIds(enable_peer_access));
|
||||
int dst_device = std::get<1>(GetDeviceIds(enable_peer_access));
|
||||
|
||||
LinearAllocGuard2D<int> src_allocation(width, height);
|
||||
HIP_CHECK(hipSetDevice(dst_device));
|
||||
LinearAllocGuard2D<int> dst_allocation(width, height);
|
||||
HIP_CHECK(hipSetDevice(src_device));
|
||||
|
||||
benchmark.Run(dst_allocation.ptr(), dst_allocation.pitch(),
|
||||
src_allocation.ptr(), src_allocation.pitch(),
|
||||
dst_allocation.width(), dst_allocation.height(),
|
||||
kind);
|
||||
}
|
||||
}
|
||||
|
||||
#if HT_NVIDIA
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyParam2D` from Device to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 4 MB x 32 B
|
||||
* - Large: 16 MB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyParam2D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyParam2D_DeviceToHost") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 32, hipMemcpyDeviceToHost);
|
||||
}
|
||||
#endif
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyParam2D` from Host to Device:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 4 MB x 32 B
|
||||
* - Large: 16 MB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyParam2D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyParam2D_HostToDevice") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 32, hipMemcpyHostToDevice);
|
||||
}
|
||||
|
||||
#if HT_NVIDIA
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyParam2D` from Host to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 4 MB x 32 B
|
||||
* - Large: 16 MB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyParam2D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyParam2D_HostToHost") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 32, hipMemcpyHostToHost);
|
||||
}
|
||||
#endif
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyParam2D` from Device to Device with peer access disabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 4 MB x 32 B
|
||||
* - Large: 16 MB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyParam2D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyParam2D_DeviceToDevice_DisablePeerAccess") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 32, hipMemcpyDeviceToDevice);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyParam2D` from Device to Device with peer access enabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 4 MB x 32 B
|
||||
* - Large: 16 MB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyParam2D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Multi-device
|
||||
* - Device supports Peer-to-Peer access
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyParam2D_DeviceToDevice_EnablePeerAccess") {
|
||||
if (HipTest::getDeviceCount() < 2) {
|
||||
HipTest::HIP_SKIP_TEST("This test requires 2 GPUs. Skipping.");
|
||||
return;
|
||||
}
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 32, hipMemcpyDeviceToDevice, true);
|
||||
}
|
||||
@@ -0,0 +1,193 @@
|
||||
/*
|
||||
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 "memcpy_performance_common.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class MemcpyParam2DBenchmark : public Benchmark<MemcpyParam2DBenchmark> {
|
||||
public:
|
||||
void operator()(void* dst, size_t dst_pitch, void* src, size_t src_pitch,
|
||||
size_t width, size_t height, hipMemcpyKind kind, const hipStream_t& stream) {
|
||||
hip_Memcpy2D params = CreateMemcpy2DParam(dst, dst_pitch, src, src_pitch,
|
||||
width, height, kind);
|
||||
TIMED_SECTION(kTimerTypeCpu) {
|
||||
HIP_CHECK(hipMemcpyParam2DAsync(¶ms, stream));
|
||||
}
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(size_t width, size_t height, hipMemcpyKind kind,
|
||||
bool enable_peer_access=false) {
|
||||
MemcpyParam2DBenchmark benchmark;
|
||||
benchmark.AddSectionName("(" + std::to_string(width) + ", " + std::to_string(height) + ")");
|
||||
|
||||
const StreamGuard stream_guard(Streams::created);
|
||||
const hipStream_t stream = stream_guard.stream();
|
||||
|
||||
if (kind == hipMemcpyDeviceToHost) {
|
||||
LinearAllocGuard2D<int> device_allocation(width, height);
|
||||
LinearAllocGuard<int> host_allocation(LinearAllocs::hipHostMalloc, device_allocation.width() * height);
|
||||
benchmark.Run(host_allocation.ptr(), device_allocation.width(),
|
||||
device_allocation.ptr(), device_allocation.pitch(),
|
||||
device_allocation.width(), device_allocation.height(),
|
||||
kind, stream);
|
||||
} else if (kind == hipMemcpyHostToDevice) {
|
||||
LinearAllocGuard2D<int> device_allocation(width, height);
|
||||
LinearAllocGuard<int> host_allocation(LinearAllocs::hipHostMalloc, device_allocation.width() * height);
|
||||
benchmark.Run(device_allocation.ptr(), device_allocation.pitch(),
|
||||
host_allocation.ptr(), device_allocation.width(),
|
||||
device_allocation.width(), device_allocation.height(),
|
||||
kind, stream);
|
||||
} else if (kind == hipMemcpyHostToHost) {
|
||||
LinearAllocGuard<int> src_allocation(LinearAllocs::hipHostMalloc, width * sizeof(int) * height);
|
||||
LinearAllocGuard<int> dst_allocation(LinearAllocs::hipHostMalloc, width * sizeof(int) * height);
|
||||
benchmark.Run(dst_allocation.ptr(), width * sizeof(int),
|
||||
src_allocation.ptr(), width * sizeof(int),
|
||||
width * sizeof(int), height, kind, stream);
|
||||
} else {
|
||||
// hipMemcpyDeviceToDevice
|
||||
int src_device = std::get<0>(GetDeviceIds(enable_peer_access));
|
||||
int dst_device = std::get<1>(GetDeviceIds(enable_peer_access));
|
||||
|
||||
LinearAllocGuard2D<int> src_allocation(width, height);
|
||||
HIP_CHECK(hipSetDevice(dst_device));
|
||||
LinearAllocGuard2D<int> dst_allocation(width, height);
|
||||
HIP_CHECK(hipSetDevice(src_device));
|
||||
benchmark.Run(dst_allocation.ptr(), dst_allocation.pitch(),
|
||||
src_allocation.ptr(), src_allocation.pitch(),
|
||||
dst_allocation.width(), dst_allocation.height(),
|
||||
kind, stream);
|
||||
}
|
||||
}
|
||||
|
||||
#if HT_NVIDIA
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyParam2DAsync` from Device to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 4 MB x 32 B
|
||||
* - Large: 16 MB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyParam2DAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyParam2DAsync_DeviceToHost") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 32, hipMemcpyDeviceToHost);
|
||||
}
|
||||
#endif
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyParam2DAsync` from Host to Device:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 4 MB x 32 B
|
||||
* - Large: 16 MB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyParam2DAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyParam2DAsync_HostToDevice") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 32, hipMemcpyHostToDevice);
|
||||
}
|
||||
|
||||
#if HT_NVIDIA
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyParam2DAsync` from Host to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 4 MB x 32 B
|
||||
* - Large: 16 MB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyParam2DAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyParam2DAsync_HostToHost") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 32, hipMemcpyHostToHost);
|
||||
}
|
||||
#endif
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyParam2DAsync` from Device to Device with peer access disabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 4 MB x 32 B
|
||||
* - Large: 16 MB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyParam2DAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyParam2DAsync_DeviceToDevice_DisablePeerAccess") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 32, hipMemcpyDeviceToDevice);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyParam2DAsync` from Device to Device with peer access enabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 4 MB x 32 B
|
||||
* - Large: 16 MB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyParam2DAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Multi-device
|
||||
* - Device supports Peer-to-Peer access
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyParam2DAsync_DeviceToDevice_EnablePeerAccess") {
|
||||
if (HipTest::getDeviceCount() < 2) {
|
||||
HipTest::HIP_SKIP_TEST("This test requires 2 GPUs. Skipping.");
|
||||
return;
|
||||
}
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 32, hipMemcpyDeviceToDevice, true);
|
||||
}
|
||||
@@ -0,0 +1,109 @@
|
||||
/*
|
||||
Copyright (c) 2024 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 "memcpy_performance_common.hh"
|
||||
#pragma clang diagnostic ignored "-Wvla-extension"
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
__device__ int devSymbol[1_MB];
|
||||
|
||||
class MemcpyToSymbolBenchmark : public Benchmark<MemcpyToSymbolBenchmark> {
|
||||
public:
|
||||
void operator()(const void* source, size_t size, size_t offset) {
|
||||
TIMED_SECTION(kTimerTypeCpu) {
|
||||
HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), source, size, offset));
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(const void* source, size_t size=1, size_t offset=0) {
|
||||
MemcpyToSymbolBenchmark benchmark;
|
||||
benchmark.AddSectionName(std::to_string(size));
|
||||
benchmark.AddSectionName(std::to_string(offset));
|
||||
benchmark.Run(source, size, offset);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyToSymbol` from Host to Device.
|
||||
* - Utilizes sigular integer values.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyToSymbol.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyToSymbol_SingularValue") {
|
||||
int set{42};
|
||||
RunBenchmark(&set);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyToSymbol` from Host to Device.
|
||||
* - Utilizes array integers:
|
||||
* - Small: 1 KB
|
||||
* - Medium: 4 KB
|
||||
* - Large: 1 MB
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyToSymbol.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyToSymbol_ArrayValue") {
|
||||
size_t size = GENERATE(1_KB, 4_KB, 1_MB);
|
||||
std::vector<int> array(size);
|
||||
std::fill_n(array.data(), size, 42);
|
||||
|
||||
RunBenchmark(array.data(), sizeof(int) * size);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyToSymbol` from Host to Device.
|
||||
* - Utilizes array integers with offsets:
|
||||
* - Small: 1 KB
|
||||
* - Medium: 4 KB
|
||||
* - Large: 1 MB
|
||||
* - Offset: 0 and size/2
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyToSymbol.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyToSymbol_WithOffset") {
|
||||
size_t size = GENERATE(1_KB, 4_KB, 1_MB);
|
||||
std::vector<int> array(size);
|
||||
std::fill_n(array.data(), size, 42);
|
||||
|
||||
size_t offset = GENERATE_REF(0, size / 2);
|
||||
RunBenchmark(array.data() + offset, sizeof(int) * (size - offset), offset * sizeof(int));
|
||||
}
|
||||
@@ -0,0 +1,116 @@
|
||||
/*
|
||||
Copyright (c) 2024 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 "memcpy_performance_common.hh"
|
||||
#pragma clang diagnostic ignored "-Wvla-extension"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
__device__ int devSymbol[1_MB];
|
||||
|
||||
class MemcpyToSymbolAsyncBenchmark : public Benchmark<MemcpyToSymbolAsyncBenchmark> {
|
||||
public:
|
||||
void operator()(const void* source, size_t size, size_t offset, const hipStream_t& stream) {
|
||||
TIMED_SECTION_STREAM(kTimerTypeEvent, stream) {
|
||||
HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), source, size, offset,
|
||||
hipMemcpyHostToDevice, stream));
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(const void* source, size_t size=1, size_t offset=0) {
|
||||
MemcpyToSymbolAsyncBenchmark benchmark;
|
||||
benchmark.AddSectionName(std::to_string(size));
|
||||
benchmark.AddSectionName(std::to_string(offset));
|
||||
|
||||
const StreamGuard stream_guard(Streams::created);
|
||||
const hipStream_t stream = stream_guard.stream();
|
||||
benchmark.Run(source, size, offset, stream);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyToSymbolAsync` from Host to Device.
|
||||
* - Utilizes sigular integer values.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyToSymbolAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyToSymbolAsync_SingularValue") {
|
||||
int set{42};
|
||||
RunBenchmark(&set);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyToSymbolAsync` from Host to Device.
|
||||
* - Utilizes array integers:
|
||||
* - Small: 1 KB
|
||||
* - Medium: 4 KB
|
||||
* - Large: 1 MB
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyToSymbolAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyToSymbolAsync_ArrayValue") {
|
||||
size_t size = GENERATE(1_KB, 4_KB, 1_MB);
|
||||
std::vector<int> array(size);
|
||||
std::fill_n(array.data(), size, 42);
|
||||
|
||||
RunBenchmark(array.data(), sizeof(int) * size);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyToSymbolAsync` from Host to Device.
|
||||
* - Utilizes array integers with offsets:
|
||||
* - Small: 1 KB
|
||||
* - Medium: 4 KB
|
||||
* - Large: 1 MB
|
||||
* - Offset: 0 and size/2
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyToSymbolAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyToSymbolAsync_WithOffset") {
|
||||
size_t size = GENERATE(1_KB, 4_KB, 1_MB);
|
||||
std::vector<int> array(size);
|
||||
std::fill_n(array.data(), size, 42);
|
||||
|
||||
size_t offset = GENERATE_REF(0, size / 2);
|
||||
RunBenchmark(array.data() + offset, sizeof(int) * (size - offset), offset * sizeof(int));
|
||||
}
|
||||
@@ -0,0 +1,192 @@
|
||||
/*
|
||||
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 "memcpy_performance_common.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup memcpy memcpy
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class MemcpyWithStreamBenchmark : public Benchmark<MemcpyWithStreamBenchmark> {
|
||||
public:
|
||||
void operator()(void* dst, const void* src, size_t size, hipMemcpyKind kind, hipStream_t stream) {
|
||||
TIMED_SECTION(kTimerTypeCpu) {
|
||||
HIP_CHECK(hipMemcpyWithStream(dst, src, size, kind, stream));
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
static void RunBenchmark(LinearAllocs dst_allocation_type, LinearAllocs src_allocation_type,
|
||||
size_t size, hipMemcpyKind kind, bool enable_peer_access=false) {
|
||||
MemcpyWithStreamBenchmark benchmark;
|
||||
benchmark.AddSectionName(std::to_string(size));
|
||||
benchmark.AddSectionName(GetAllocationSectionName(src_allocation_type));
|
||||
benchmark.AddSectionName(GetAllocationSectionName(dst_allocation_type));
|
||||
|
||||
const StreamGuard stream_guard(Streams::created);
|
||||
const hipStream_t stream = stream_guard.stream();
|
||||
|
||||
if (kind != hipMemcpyDeviceToDevice) {
|
||||
LinearAllocGuard<int> src_allocation(src_allocation_type, size);
|
||||
LinearAllocGuard<int> dst_allocation(dst_allocation_type, size);
|
||||
benchmark.Run(dst_allocation.ptr(), src_allocation.ptr(), size, kind, stream);
|
||||
} else {
|
||||
int src_device = std::get<0>(GetDeviceIds(enable_peer_access));
|
||||
int dst_device = std::get<1>(GetDeviceIds(enable_peer_access));
|
||||
|
||||
LinearAllocGuard<int> src_allocation(LinearAllocs::hipMalloc, size);
|
||||
HIP_CHECK(hipSetDevice(dst_device));
|
||||
LinearAllocGuard<int> dst_allocation(LinearAllocs::hipMalloc, size);
|
||||
HIP_CHECK(hipSetDevice(src_device));
|
||||
benchmark.Run(dst_allocation.ptr(), src_allocation.ptr(), size, kind, stream);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyWithStream` from Device to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: device malloc
|
||||
* - Destination: host pinned and pageable
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyWithStream.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyWithStream_DeviceToHost") {
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto src_allocation_type = LinearAllocs::hipMalloc;
|
||||
const auto dst_allocation_type = GENERATE(LinearAllocs::malloc, LinearAllocs::hipHostMalloc);
|
||||
RunBenchmark(dst_allocation_type, src_allocation_type, allocation_size, hipMemcpyDeviceToHost);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyWithStream` from Host to Device:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: host pinned and pageable
|
||||
* - Destination: device malloc
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyWithStream.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyWithStream_HostToDevice") {
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto src_allocation_type = GENERATE(LinearAllocs::malloc, LinearAllocs::hipHostMalloc);
|
||||
const auto dst_allocation_type = LinearAllocs::hipMalloc;
|
||||
RunBenchmark(dst_allocation_type, src_allocation_type, allocation_size, hipMemcpyHostToDevice);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyWithStream` from Host to Host:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: host pinned and pageable
|
||||
* - Destination: host pinned and pageable
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyWithStream.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyWithStream_HostToHost") {
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto src_allocation_type = GENERATE(LinearAllocs::malloc, LinearAllocs::hipHostMalloc);
|
||||
const auto dst_allocation_type = GENERATE(LinearAllocs::malloc, LinearAllocs::hipHostMalloc);
|
||||
RunBenchmark(dst_allocation_type, src_allocation_type, allocation_size, hipMemcpyHostToHost);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpy` from Device to Device with peer access disabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: device malloc
|
||||
* - Destination: device malloc
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyWithStream.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyWithStream_DeviceToDevice_DisablePeerAccess") {
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto src_allocation_type = LinearAllocs::hipMalloc;
|
||||
const auto dst_allocation_type = LinearAllocs::hipMalloc;
|
||||
RunBenchmark(dst_allocation_type, src_allocation_type, allocation_size, hipMemcpyDeviceToDevice);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemcpyWithStream` from Device to Device with peer access enabled:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - Source: device malloc
|
||||
* - Destination: device malloc
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memcpy/hipMemcpyWithStream.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Multi-device
|
||||
* - Device supports Peer-to-Peer access
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemcpyWithStream_DeviceToDevice_EnablePeerAccess") {
|
||||
if (HipTest::getDeviceCount() < 2) {
|
||||
HipTest::HIP_SKIP_TEST("This test requires 2 GPUs. Skipping.");
|
||||
return;
|
||||
}
|
||||
const auto allocation_size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto src_allocation_type = LinearAllocs::hipMalloc;
|
||||
const auto dst_allocation_type = LinearAllocs::hipMalloc;
|
||||
RunBenchmark(dst_allocation_type, src_allocation_type, allocation_size, hipMemcpyDeviceToDevice, true);
|
||||
}
|
||||
@@ -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.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <performance_common.hh>
|
||||
|
||||
static hip_Memcpy2D CreateMemcpy2DParam(void* dst, size_t dpitch, void* src, size_t spitch,
|
||||
size_t width, size_t height, hipMemcpyKind kind) {
|
||||
hip_Memcpy2D params = {};
|
||||
memset(¶ms, 0, sizeof(hip_Memcpy2D));
|
||||
|
||||
const hipExtent src_offset = {0, 0, 0};
|
||||
const hipExtent dst_offset = {0, 0, 0};
|
||||
|
||||
params.dstPitch = dpitch;
|
||||
switch (kind) {
|
||||
case hipMemcpyDeviceToHost:
|
||||
case hipMemcpyHostToHost:
|
||||
params.dstMemoryType = hipMemoryTypeHost;
|
||||
params.dstHost = dst;
|
||||
break;
|
||||
case hipMemcpyDeviceToDevice:
|
||||
case hipMemcpyHostToDevice:
|
||||
params.dstMemoryType = hipMemoryTypeDevice;
|
||||
params.dstDevice = reinterpret_cast<hipDeviceptr_t>(dst);
|
||||
break;
|
||||
default:
|
||||
REQUIRE(false);
|
||||
}
|
||||
|
||||
params.srcPitch = dpitch;
|
||||
switch (kind) {
|
||||
case hipMemcpyDeviceToHost:
|
||||
case hipMemcpyHostToHost:
|
||||
params.srcMemoryType = hipMemoryTypeHost;
|
||||
params.srcHost = src;
|
||||
break;
|
||||
case hipMemcpyDeviceToDevice:
|
||||
case hipMemcpyHostToDevice:
|
||||
params.srcMemoryType = hipMemoryTypeDevice;
|
||||
params.srcDevice = reinterpret_cast<hipDeviceptr_t>(src);
|
||||
break;
|
||||
default:
|
||||
REQUIRE(false);
|
||||
}
|
||||
|
||||
params.WidthInBytes = width;
|
||||
params.Height = height;
|
||||
params.srcXInBytes = src_offset.width;
|
||||
params.srcY = src_offset.height;
|
||||
params.dstXInBytes = dst_offset.width;
|
||||
params.dstY = dst_offset.height;
|
||||
|
||||
return params;
|
||||
}
|
||||
|
||||
static hipMemcpy3DParms CreateMemcpy3DParam(hipPitchedPtr dst_ptr, hipPos dst_pos,
|
||||
hipPitchedPtr src_ptr, hipPos src_pos,
|
||||
hipExtent extent, hipMemcpyKind kind) {
|
||||
hipMemcpy3DParms params = {};
|
||||
memset(¶ms, 0, sizeof(hipMemcpy3DParms));
|
||||
params.dstPtr = dst_ptr;
|
||||
params.dstPos = dst_pos;
|
||||
params.srcPtr = src_ptr;
|
||||
params.srcPos = src_pos;
|
||||
params.extent = extent;
|
||||
params.kind = kind;
|
||||
return params;
|
||||
}
|
||||
|
||||
static std::tuple<int, int> GetDeviceIds(bool enable_peer_access) {
|
||||
int src_device = 0;
|
||||
int dst_device = 1;
|
||||
|
||||
if (enable_peer_access) {
|
||||
int can_access_peer = 0;
|
||||
HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, src_device, dst_device));
|
||||
if (!can_access_peer) {
|
||||
INFO("Peer access cannot be enabled between devices " << src_device << " and " << dst_device);
|
||||
REQUIRE(can_access_peer);
|
||||
}
|
||||
HIP_CHECK(hipDeviceEnablePeerAccess(dst_device, 0));
|
||||
} else {
|
||||
dst_device = 0;
|
||||
}
|
||||
|
||||
return {src_device, dst_device};
|
||||
}
|
||||
@@ -0,0 +1,39 @@
|
||||
# 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.
|
||||
|
||||
set(TEST_SRC
|
||||
hipMemset.cc
|
||||
hipMemsetAsync.cc
|
||||
hipMemsetD8.cc
|
||||
hipMemsetD8Async.cc
|
||||
hipMemsetD16.cc
|
||||
hipMemsetD16Async.cc
|
||||
hipMemsetD32.cc
|
||||
hipMemsetD32Async.cc
|
||||
hipMemset2D.cc
|
||||
hipMemset2DAsync.cc
|
||||
hipMemset3D.cc
|
||||
hipMemset3DAsync.cc
|
||||
)
|
||||
|
||||
hip_add_exe_to_target(NAME MemsetPerformance
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME build_tests
|
||||
COMPILE_OPTIONS -std=c++17)
|
||||
@@ -0,0 +1,79 @@
|
||||
/*
|
||||
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 <hip_test_common.hh>
|
||||
#include <performance_common.hh>
|
||||
#include <resource_guards.hh>
|
||||
|
||||
/**
|
||||
* @addtogroup memset memset
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
* Contains performance tests for all memset HIP APIs.
|
||||
*/
|
||||
|
||||
class MemsetBenchmark : public Benchmark<MemsetBenchmark> {
|
||||
public:
|
||||
MemsetBenchmark(LinearAllocs allocation_type, size_t size)
|
||||
: dst_(allocation_type, size), size_(size) {}
|
||||
|
||||
void operator()() {
|
||||
TIMED_SECTION(kTimerTypeEvent) { HIP_CHECK(hipMemset(dst_.ptr(), 17, size_)); }
|
||||
}
|
||||
|
||||
private:
|
||||
LinearAllocGuard<void> dst_;
|
||||
const size_t size_;
|
||||
};
|
||||
|
||||
static void RunBenchmark(LinearAllocs allocation_type, size_t size) {
|
||||
MemsetBenchmark benchmark(allocation_type, size);
|
||||
benchmark.AddSectionName(std::to_string(size));
|
||||
benchmark.AddSectionName(GetAllocationSectionName(allocation_type));
|
||||
benchmark.Run();
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemset`:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - device
|
||||
* - host
|
||||
* - managed
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memset/hipMemset.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemset") {
|
||||
const auto size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto allocation_type = GENERATE(LinearAllocs::hipMalloc, LinearAllocs::hipHostMalloc,
|
||||
LinearAllocs::hipMallocManaged);
|
||||
RunBenchmark(allocation_type, size);
|
||||
}
|
||||
@@ -0,0 +1,71 @@
|
||||
/*
|
||||
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 <hip_test_common.hh>
|
||||
#include <performance_common.hh>
|
||||
#include <resource_guards.hh>
|
||||
|
||||
/**
|
||||
* @addtogroup memset memset
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class Memset2DBenchmark : public Benchmark<Memset2DBenchmark> {
|
||||
public:
|
||||
Memset2DBenchmark(size_t width, size_t height) : dst_(width, height) {}
|
||||
|
||||
void operator()() {
|
||||
TIMED_SECTION(kTimerTypeEvent) {
|
||||
HIP_CHECK(hipMemset2D(dst_.ptr(), dst_.pitch(), 17, dst_.width(), dst_.height()));
|
||||
}
|
||||
}
|
||||
|
||||
private:
|
||||
LinearAllocGuard2D<char> dst_;
|
||||
};
|
||||
|
||||
static void RunBenchmark(size_t width, size_t height) {
|
||||
Memset2DBenchmark benchmark(width, height);
|
||||
benchmark.AddSectionName("(" + std::to_string(width) + ", " + std::to_string(height) + ")");
|
||||
benchmark.Run();
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemset2D`:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 4 MB x 32 B
|
||||
* - Large: 16 MB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memset/hipMemset2D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemset2D") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 32);
|
||||
}
|
||||
@@ -0,0 +1,74 @@
|
||||
/*
|
||||
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 <hip_test_common.hh>
|
||||
#include <performance_common.hh>
|
||||
#include <resource_guards.hh>
|
||||
|
||||
/**
|
||||
* @addtogroup memset memset
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class Memset2DAsyncBenchmark : public Benchmark<Memset2DAsyncBenchmark> {
|
||||
public:
|
||||
Memset2DAsyncBenchmark(size_t width, size_t height)
|
||||
: dst_(width, height), stream_(Streams::created) {}
|
||||
|
||||
void operator()(size_t width, size_t height) {
|
||||
TIMED_SECTION_STREAM(kTimerTypeEvent, stream_.stream()) {
|
||||
HIP_CHECK(hipMemset2DAsync(dst_.ptr(), dst_.pitch(), 17, dst_.width(), dst_.height(),
|
||||
stream_.stream()));
|
||||
}
|
||||
}
|
||||
|
||||
private:
|
||||
LinearAllocGuard2D<char> dst_;
|
||||
StreamGuard stream_;
|
||||
};
|
||||
|
||||
static void RunBenchmark(size_t width, size_t height) {
|
||||
Memset2DAsyncBenchmark benchmark(width, height);
|
||||
benchmark.AddSectionName("(" + std::to_string(width) + ", " + std::to_string(height) + ")");
|
||||
benchmark.Run(width, height);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemset2DAsync`:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 32 B
|
||||
* - Medium: 4 MB x 32 B
|
||||
* - Large: 16 MB x 32 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memset/hipMemset2DAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemset2DAsync") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 32);
|
||||
}
|
||||
@@ -0,0 +1,72 @@
|
||||
/*
|
||||
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 <hip_test_common.hh>
|
||||
#include <performance_common.hh>
|
||||
#include <resource_guards.hh>
|
||||
|
||||
/**
|
||||
* @addtogroup memset memset
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class Memset3DBenchmark : public Benchmark<Memset3DBenchmark> {
|
||||
public:
|
||||
Memset3DBenchmark(size_t width, size_t height, size_t depth) : dst_(width, height, depth) {}
|
||||
|
||||
void operator()() {
|
||||
TIMED_SECTION(kTimerTypeEvent) {
|
||||
HIP_CHECK(hipMemset3D(dst_.pitched_ptr(), 17, dst_.extent()));
|
||||
}
|
||||
}
|
||||
|
||||
private:
|
||||
LinearAllocGuard3D<char> dst_;
|
||||
};
|
||||
|
||||
static void RunBenchmark(size_t width, size_t height, size_t depth) {
|
||||
Memset3DBenchmark benchmark(width, height, depth);
|
||||
benchmark.AddSectionName("(" + std::to_string(width) + ", " + std::to_string(height) + ", " +
|
||||
std::to_string(depth) + ")");
|
||||
benchmark.Run();
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemset3D`:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 16 B x 4 B
|
||||
* - Medium: 4 MB x 16 B x 4 B
|
||||
* - Large: 16 MB x 16 B x 4 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memset/hipMemset3D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemset3D") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 16, 4);
|
||||
}
|
||||
@@ -0,0 +1,74 @@
|
||||
/*
|
||||
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 <hip_test_common.hh>
|
||||
#include <performance_common.hh>
|
||||
#include <resource_guards.hh>
|
||||
|
||||
/**
|
||||
* @addtogroup memset memset
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class Memset3DAsyncBenchmark : public Benchmark<Memset3DAsyncBenchmark> {
|
||||
public:
|
||||
Memset3DAsyncBenchmark(size_t width, size_t height, size_t depth)
|
||||
: dst_(width, height, depth), stream_(Streams::created) {}
|
||||
|
||||
void operator()() {
|
||||
TIMED_SECTION_STREAM(kTimerTypeEvent, stream_.stream()) {
|
||||
HIP_CHECK(hipMemset3DAsync(dst_.pitched_ptr(), 17, dst_.extent(), stream_.stream()));
|
||||
}
|
||||
}
|
||||
|
||||
private:
|
||||
LinearAllocGuard3D<char> dst_;
|
||||
StreamGuard stream_;
|
||||
};
|
||||
|
||||
static void RunBenchmark(size_t width, size_t height, size_t depth) {
|
||||
Memset3DAsyncBenchmark benchmark(width, height, depth);
|
||||
benchmark.AddSectionName("(" + std::to_string(width) + ", " + std::to_string(height) + ", " +
|
||||
std::to_string(depth) + ")");
|
||||
benchmark.Run();
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemset3DAsync`:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB x 16 B x 4 B
|
||||
* - Medium: 4 MB x 16 B x 4 B
|
||||
* - Large: 16 MB x 16 B x 4 B
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memset/hipMemset3DAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemset3DAsync") {
|
||||
const auto width = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
RunBenchmark(width, 16, 4);
|
||||
}
|
||||
@@ -0,0 +1,81 @@
|
||||
/*
|
||||
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 <hip_test_common.hh>
|
||||
#include <performance_common.hh>
|
||||
#include <resource_guards.hh>
|
||||
|
||||
/**
|
||||
* @addtogroup memset memset
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class MemsetAsyncBenchmark : public Benchmark<MemsetAsyncBenchmark> {
|
||||
public:
|
||||
MemsetAsyncBenchmark(LinearAllocs allocation_type, size_t size)
|
||||
: dst_(allocation_type, size), size_(size), stream_(Streams::created) {}
|
||||
|
||||
void operator()() {
|
||||
TIMED_SECTION_STREAM(kTimerTypeEvent, stream_.stream()) {
|
||||
HIP_CHECK(hipMemsetAsync(dst_.ptr(), 17, size_, stream_.stream()));
|
||||
}
|
||||
}
|
||||
|
||||
private:
|
||||
LinearAllocGuard<void> dst_;
|
||||
const size_t size_;
|
||||
StreamGuard stream_;
|
||||
};
|
||||
|
||||
static void RunBenchmark(LinearAllocs allocation_type, size_t size) {
|
||||
MemsetAsyncBenchmark benchmark(allocation_type, size);
|
||||
benchmark.AddSectionName(std::to_string(size));
|
||||
benchmark.AddSectionName(GetAllocationSectionName(allocation_type));
|
||||
benchmark.Run();
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemsetAsync`:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - device
|
||||
* - host
|
||||
* - managed
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memset/hipMemsetAsync.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemsetAsync") {
|
||||
const auto size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto allocation_type = GENERATE(LinearAllocs::hipMalloc, LinearAllocs::hipHostMalloc,
|
||||
LinearAllocs::hipMallocManaged);
|
||||
RunBenchmark(allocation_type, size);
|
||||
}
|
||||
@@ -0,0 +1,80 @@
|
||||
/*
|
||||
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 <hip_test_common.hh>
|
||||
#include <performance_common.hh>
|
||||
#include <resource_guards.hh>
|
||||
|
||||
/**
|
||||
* @addtogroup memset memset
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class MemsetD16Benchmark : public Benchmark<MemsetD16Benchmark> {
|
||||
public:
|
||||
MemsetD16Benchmark(LinearAllocs allocation_type, size_t size)
|
||||
: dst_(allocation_type, size), size_(size) {}
|
||||
|
||||
void operator()() {
|
||||
TIMED_SECTION(kTimerTypeEvent) {
|
||||
HIP_CHECK(hipMemsetD16(reinterpret_cast<hipDeviceptr_t>(dst_.ptr()), 311, size_));
|
||||
}
|
||||
}
|
||||
|
||||
private:
|
||||
LinearAllocGuard<void> dst_;
|
||||
const size_t size_;
|
||||
};
|
||||
|
||||
static void RunBenchmark(LinearAllocs allocation_type, size_t size) {
|
||||
MemsetD16Benchmark benchmark(allocation_type, size);
|
||||
benchmark.AddSectionName(std::to_string(size));
|
||||
benchmark.AddSectionName(GetAllocationSectionName(allocation_type));
|
||||
benchmark.Run();
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemsetD16`:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - device
|
||||
* - host
|
||||
* - managed
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memset/hipMemsetD16.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemsetD16") {
|
||||
const auto size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto allocation_type = GENERATE(LinearAllocs::hipMalloc, LinearAllocs::hipHostMalloc,
|
||||
LinearAllocs::hipMallocManaged);
|
||||
RunBenchmark(allocation_type, size);
|
||||
}
|
||||
@@ -0,0 +1,82 @@
|
||||
/*
|
||||
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 <hip_test_common.hh>
|
||||
#include <performance_common.hh>
|
||||
#include <resource_guards.hh>
|
||||
|
||||
/**
|
||||
* @addtogroup memset memset
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class MemsetD16AsyncBenchmark : public Benchmark<MemsetD16AsyncBenchmark> {
|
||||
public:
|
||||
MemsetD16AsyncBenchmark(LinearAllocs allocation_type, size_t size)
|
||||
: dst_(allocation_type, size), size_(size), stream_(Streams::created) {}
|
||||
|
||||
void operator()() {
|
||||
TIMED_SECTION_STREAM(kTimerTypeEvent, stream_.stream()) {
|
||||
HIP_CHECK(hipMemsetD16Async(reinterpret_cast<hipDeviceptr_t>(dst_.ptr()), 311, size_,
|
||||
stream_.stream()));
|
||||
}
|
||||
}
|
||||
|
||||
private:
|
||||
LinearAllocGuard<void> dst_;
|
||||
const size_t size_;
|
||||
StreamGuard stream_;
|
||||
};
|
||||
|
||||
static void RunBenchmark(LinearAllocs allocation_type, size_t size) {
|
||||
MemsetD16AsyncBenchmark benchmark(allocation_type, size);
|
||||
benchmark.AddSectionName(std::to_string(size));
|
||||
benchmark.AddSectionName(GetAllocationSectionName(allocation_type));
|
||||
benchmark.Run();
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemsetD16Async`:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - device
|
||||
* - host
|
||||
* - managed
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memset/hipMemsetD16Async.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemsetD16Async") {
|
||||
const auto size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto allocation_type = GENERATE(LinearAllocs::hipMalloc, LinearAllocs::hipHostMalloc,
|
||||
LinearAllocs::hipMallocManaged);
|
||||
RunBenchmark(allocation_type, size);
|
||||
}
|
||||
@@ -0,0 +1,80 @@
|
||||
/*
|
||||
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 <hip_test_common.hh>
|
||||
#include <performance_common.hh>
|
||||
#include <resource_guards.hh>
|
||||
|
||||
/**
|
||||
* @addtogroup memset memset
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class MemsetD32Benchmark : public Benchmark<MemsetD32Benchmark> {
|
||||
public:
|
||||
MemsetD32Benchmark(LinearAllocs allocation_type, size_t size)
|
||||
: dst_(allocation_type, size), size_(size) {}
|
||||
|
||||
void operator()() {
|
||||
TIMED_SECTION(kTimerTypeEvent) {
|
||||
HIP_CHECK(hipMemsetD32(reinterpret_cast<hipDeviceptr_t>(dst_.ptr()), 123'456, size_));
|
||||
}
|
||||
}
|
||||
|
||||
private:
|
||||
LinearAllocGuard<void> dst_;
|
||||
const size_t size_;
|
||||
};
|
||||
|
||||
static void RunBenchmark(LinearAllocs allocation_type, size_t size) {
|
||||
MemsetD32Benchmark benchmark(allocation_type, size);
|
||||
benchmark.AddSectionName(std::to_string(size));
|
||||
benchmark.AddSectionName(GetAllocationSectionName(allocation_type));
|
||||
benchmark.Run();
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemsetD32`:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - device
|
||||
* - host
|
||||
* - managed
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memset/hipMemsetD32.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemsetD32") {
|
||||
const auto size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto allocation_type = GENERATE(LinearAllocs::hipMalloc, LinearAllocs::hipHostMalloc,
|
||||
LinearAllocs::hipMallocManaged);
|
||||
RunBenchmark(allocation_type, size);
|
||||
}
|
||||
@@ -0,0 +1,82 @@
|
||||
/*
|
||||
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 <hip_test_common.hh>
|
||||
#include <performance_common.hh>
|
||||
#include <resource_guards.hh>
|
||||
|
||||
/**
|
||||
* @addtogroup memset memset
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class MemsetD32AsyncBenchmark : public Benchmark<MemsetD32AsyncBenchmark> {
|
||||
public:
|
||||
MemsetD32AsyncBenchmark(LinearAllocs allocation_type, size_t size)
|
||||
: dst_(allocation_type, size), size_(size), stream_(Streams::created) {}
|
||||
|
||||
void operator()() {
|
||||
TIMED_SECTION_STREAM(kTimerTypeEvent, stream_.stream()) {
|
||||
HIP_CHECK(hipMemsetD32Async(reinterpret_cast<hipDeviceptr_t>(dst_.ptr()), 123'456, size_,
|
||||
stream_.stream()));
|
||||
}
|
||||
}
|
||||
|
||||
private:
|
||||
LinearAllocGuard<void> dst_;
|
||||
const size_t size_;
|
||||
StreamGuard stream_;
|
||||
};
|
||||
|
||||
static void RunBenchmark(LinearAllocs allocation_type, size_t size) {
|
||||
MemsetD32AsyncBenchmark benchmark(allocation_type, size);
|
||||
benchmark.AddSectionName(std::to_string(size));
|
||||
benchmark.AddSectionName(GetAllocationSectionName(allocation_type));
|
||||
benchmark.Run();
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemsetD32Async`:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - device
|
||||
* - host
|
||||
* - managed
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memset/hipMemsetD32Async.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemsetD32Async") {
|
||||
const auto size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto allocation_type = GENERATE(LinearAllocs::hipMalloc, LinearAllocs::hipHostMalloc,
|
||||
LinearAllocs::hipMallocManaged);
|
||||
RunBenchmark(allocation_type, size);
|
||||
}
|
||||
@@ -0,0 +1,80 @@
|
||||
/*
|
||||
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 <hip_test_common.hh>
|
||||
#include <performance_common.hh>
|
||||
#include <resource_guards.hh>
|
||||
|
||||
/**
|
||||
* @addtogroup memset memset
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class MemsetD8Benchmark : public Benchmark<MemsetD8Benchmark> {
|
||||
public:
|
||||
MemsetD8Benchmark(LinearAllocs allocation_type, size_t size)
|
||||
: dst_(allocation_type, size), size_(size) {}
|
||||
|
||||
void operator()() {
|
||||
TIMED_SECTION(kTimerTypeEvent) {
|
||||
HIP_CHECK(hipMemsetD8(reinterpret_cast<hipDeviceptr_t>(dst_.ptr()), 17, size_));
|
||||
}
|
||||
}
|
||||
|
||||
private:
|
||||
LinearAllocGuard<void> dst_;
|
||||
const size_t size_;
|
||||
};
|
||||
|
||||
static void RunBenchmark(LinearAllocs allocation_type, size_t size) {
|
||||
MemsetD8Benchmark benchmark(allocation_type, size);
|
||||
benchmark.AddSectionName(std::to_string(size));
|
||||
benchmark.AddSectionName(GetAllocationSectionName(allocation_type));
|
||||
benchmark.Run();
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemsetD8`:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - device
|
||||
* - host
|
||||
* - managed
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memset/hipMemsetD8.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemsetD8") {
|
||||
const auto size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto allocation_type = GENERATE(LinearAllocs::hipMalloc, LinearAllocs::hipHostMalloc,
|
||||
LinearAllocs::hipMallocManaged);
|
||||
RunBenchmark(allocation_type, size);
|
||||
}
|
||||
@@ -0,0 +1,82 @@
|
||||
/*
|
||||
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 <hip_test_common.hh>
|
||||
#include <performance_common.hh>
|
||||
#include <resource_guards.hh>
|
||||
|
||||
/**
|
||||
* @addtogroup memset memset
|
||||
* @{
|
||||
* @ingroup PerformanceTest
|
||||
*/
|
||||
|
||||
class MemsetD8AsyncBenchmark : public Benchmark<MemsetD8AsyncBenchmark> {
|
||||
public:
|
||||
MemsetD8AsyncBenchmark(LinearAllocs allocation_type, size_t size)
|
||||
: dst_(allocation_type, size), size_(size), stream_(Streams::created) {}
|
||||
|
||||
void operator()() {
|
||||
TIMED_SECTION_STREAM(kTimerTypeEvent, stream_.stream()) {
|
||||
HIP_CHECK(hipMemsetD8Async(reinterpret_cast<hipDeviceptr_t>(dst_.ptr()), 17, size_,
|
||||
stream_.stream()));
|
||||
}
|
||||
}
|
||||
|
||||
private:
|
||||
LinearAllocGuard<void> dst_;
|
||||
const size_t size_;
|
||||
StreamGuard stream_;
|
||||
};
|
||||
|
||||
static void RunBenchmark(LinearAllocs allocation_type, size_t size) {
|
||||
MemsetD8AsyncBenchmark benchmark(allocation_type, size);
|
||||
benchmark.AddSectionName(std::to_string(size));
|
||||
benchmark.AddSectionName(GetAllocationSectionName(allocation_type));
|
||||
benchmark.Run();
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Executes `hipMemsetD8Async`:
|
||||
* -# Allocation size
|
||||
* - Small: 4 KB
|
||||
* - Medium: 4 MB
|
||||
* - Large: 16 MB
|
||||
* -# Allocation type
|
||||
* - device
|
||||
* - host
|
||||
* - managed
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - performance/memset/hipMemsetD8Async.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Performance_hipMemsetD8Async") {
|
||||
const auto size = GENERATE(4_KB, 4_MB, 16_MB);
|
||||
const auto allocation_type = GENERATE(LinearAllocs::hipMalloc, LinearAllocs::hipHostMalloc,
|
||||
LinearAllocs::hipMallocManaged);
|
||||
RunBenchmark(allocation_type, size);
|
||||
}
|
||||
Reference in New Issue
Block a user