SWDEV-403471 - [catch2][dtest] Converting perftests-memory files from HIT to catch2 (#342)
Change-Id: I13d2513f31dffe0b280039c888a97cc0d7bba31f
[ROCm/hip-tests commit: cf174d5a47]
このコミットが含まれているのは:
ROCm CI Service Account
GitHub
@@ -241,6 +241,7 @@ add_subdirectory(kernels ${CATCH_BUILD_DIR}/kernels)
|
||||
add_subdirectory(hipTestMain ${CATCH_BUILD_DIR}/hipTestMain)
|
||||
add_subdirectory(stress ${CATCH_BUILD_DIR}/stress)
|
||||
add_subdirectory(TypeQualifiers ${CATCH_BUILD_DIR}/TypeQualifiers)
|
||||
add_subdirectory(perftests ${CATCH_BUILD_DIR}/perftests)
|
||||
if(UNIX)
|
||||
add_subdirectory(multiproc ${CATCH_BUILD_DIR}/multiproc)
|
||||
endif()
|
||||
|
||||
@@ -0,0 +1,24 @@
|
||||
# Copyright (c) 2023 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.
|
||||
|
||||
add_custom_target(perf_test COMMAND "${CMAKE_CTEST_COMMAND}" -R "Perf_"
|
||||
COMMENT "Build complete, now executing the performnce test ...")
|
||||
|
||||
add_subdirectory(memory)
|
||||
@@ -0,0 +1,47 @@
|
||||
# Copyright (c) 2023 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
|
||||
hipPerfMemcpy.cc
|
||||
hipPerfBufferCopyRectSpeed.cc
|
||||
hipPerfBufferCopySpeed.cc
|
||||
hipPerfDevMemReadSpeed.cc
|
||||
hipPerfDevMemWriteSpeed.cc
|
||||
hipPerfMemFill.cc
|
||||
hipPerfMemMallocCpyFree.cc
|
||||
hipPerfMemset.cc
|
||||
hipPerfSampleRate.cc
|
||||
hipPerfSharedMemReadSpeed.cc
|
||||
)
|
||||
|
||||
if(HIP_PLATFORM MATCHES "amd")
|
||||
set(TEST_SRC
|
||||
${TEST_SRC}
|
||||
hipPerfHostNumaAlloc.cc)
|
||||
|
||||
hip_add_exe_to_target(NAME perfMemoryTest
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME perf_test
|
||||
LINKER_LIBS numa)
|
||||
else()
|
||||
hip_add_exe_to_target(NAME perfMemoryTest
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME perf_test)
|
||||
endif()
|
||||
@@ -0,0 +1,242 @@
|
||||
/*
|
||||
Copyright (c) 2023 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.
|
||||
*/
|
||||
|
||||
/**
|
||||
* @addtogroup hipMemcpy2DAsync hipMemcpy2DAsync
|
||||
* @{
|
||||
* @ingroup perfMemoryTest
|
||||
* `hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch,
|
||||
* size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream = 0)` -
|
||||
* Copies data between host and device.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
#define NUM_SIZES 8
|
||||
// 4KB, 8KB, 64KB, 256KB, 1 MB, 4MB, 16 MB, 16MB+10
|
||||
static const unsigned int Sizes[NUM_SIZES] =
|
||||
{4096, 8192, 65536, 262144, 1048576, 4194304, 16777216, 16777216+10};
|
||||
|
||||
static const unsigned int Iterations[2] = {1, 1000};
|
||||
|
||||
#define BUF_TYPES 4
|
||||
// 16 ways to combine 4 different buffer types
|
||||
#define NUM_SUBTESTS (BUF_TYPES*BUF_TYPES)
|
||||
|
||||
static void setData(void *ptr, unsigned int size, char value) {
|
||||
char *ptr2 = reinterpret_cast<char *>(ptr);
|
||||
for (unsigned int i = 0; i < size ; i++) {
|
||||
ptr2[i] = value;
|
||||
}
|
||||
}
|
||||
|
||||
static bool hipPerfBufferCopyRectSpeed_test(int p_tests) {
|
||||
unsigned int bufSize_;
|
||||
unsigned int numIter;
|
||||
bool hostMalloc[2] = {false};
|
||||
bool hostRegister[2] = {false};
|
||||
bool unpinnedMalloc[2] = {false};
|
||||
void *memptr[2] = {NULL};
|
||||
void *alignedmemptr[2] = {NULL};
|
||||
void *srcBuffer = NULL;
|
||||
void *dstBuffer = NULL;
|
||||
|
||||
int numTests = (p_tests == -1) ? (NUM_SIZES*NUM_SUBTESTS*2 - 1) : p_tests;
|
||||
int test = (p_tests == -1) ? 0 : p_tests;
|
||||
|
||||
for ( ; test <= numTests ; test++ ) {
|
||||
unsigned int srcTest = (test / NUM_SIZES) % BUF_TYPES;
|
||||
unsigned int dstTest = (test / (NUM_SIZES*BUF_TYPES)) % BUF_TYPES;
|
||||
bufSize_ = Sizes[test % NUM_SIZES];
|
||||
hostMalloc[0] = hostMalloc[1] = false;
|
||||
hostRegister[0] = hostRegister[1] = false;
|
||||
unpinnedMalloc[0] = unpinnedMalloc[1] = false;
|
||||
srcBuffer = dstBuffer = 0;
|
||||
memptr[0] = memptr[1] = 0;
|
||||
alignedmemptr[0] = alignedmemptr[1] = NULL;
|
||||
|
||||
size_t width = static_cast<size_t>(sqrt(static_cast<float>(bufSize_)));
|
||||
|
||||
if (srcTest == 3) {
|
||||
hostRegister[0] = true;
|
||||
} else if (srcTest == 2) {
|
||||
hostMalloc[0] = true;
|
||||
} else if (srcTest == 1) {
|
||||
unpinnedMalloc[0] = true;
|
||||
}
|
||||
|
||||
if (dstTest == 1) {
|
||||
unpinnedMalloc[1] = true;
|
||||
} else if (dstTest == 2) {
|
||||
hostMalloc[1] = true;
|
||||
} else if (dstTest == 3) {
|
||||
hostRegister[1] = true;
|
||||
}
|
||||
|
||||
numIter = Iterations[test / (NUM_SIZES * NUM_SUBTESTS)];
|
||||
|
||||
if (hostMalloc[0]) {
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&srcBuffer),
|
||||
bufSize_, 0));
|
||||
setData(srcBuffer, bufSize_, 0xd0);
|
||||
} else if (hostRegister[0]) {
|
||||
memptr[0] = malloc(bufSize_ + 4096);
|
||||
alignedmemptr[0] = reinterpret_cast<void*>(memptr[0]);
|
||||
srcBuffer = alignedmemptr[0];
|
||||
setData(srcBuffer, bufSize_, 0xd0);
|
||||
HIP_CHECK(hipHostRegister(srcBuffer, bufSize_, 0));
|
||||
} else if (unpinnedMalloc[0]) {
|
||||
memptr[0] = malloc(bufSize_ + 4096);
|
||||
alignedmemptr[0] = reinterpret_cast<void*>(memptr[0]);
|
||||
srcBuffer = alignedmemptr[0];
|
||||
setData(srcBuffer, bufSize_, 0xd0);
|
||||
} else {
|
||||
HIP_CHECK(hipMalloc(&srcBuffer, bufSize_));
|
||||
HIP_CHECK(hipMemset(srcBuffer, 0xd0, bufSize_));
|
||||
}
|
||||
|
||||
if (hostMalloc[1]) {
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&dstBuffer),
|
||||
bufSize_, 0));
|
||||
} else if (hostRegister[1]) {
|
||||
memptr[1] = malloc(bufSize_ + 4096);
|
||||
alignedmemptr[1] = reinterpret_cast<void*>(memptr[0]);
|
||||
dstBuffer = alignedmemptr[1];
|
||||
HIP_CHECK(hipHostRegister(dstBuffer, bufSize_, 0));
|
||||
} else if (unpinnedMalloc[1]) {
|
||||
memptr[1] = malloc(bufSize_ + 4096);
|
||||
alignedmemptr[1] = reinterpret_cast<void*>(memptr[0]);
|
||||
dstBuffer = alignedmemptr[1];
|
||||
} else {
|
||||
HIP_CHECK(hipMalloc(&dstBuffer, bufSize_));
|
||||
}
|
||||
|
||||
// warm up
|
||||
HIP_CHECK(hipMemcpy2D(dstBuffer, width, srcBuffer,
|
||||
width, width, width, hipMemcpyDefault));
|
||||
|
||||
// measure performance based on host time
|
||||
auto all_start = std::chrono::steady_clock::now();
|
||||
|
||||
for (unsigned int i = 0; i < numIter; i++) {
|
||||
HIP_CHECK(hipMemcpy2DAsync(dstBuffer, width, srcBuffer,
|
||||
width, width, width, hipMemcpyDefault, NULL));
|
||||
}
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
auto all_end = std::chrono::steady_clock::now();
|
||||
std::chrono::duration<double> elapsed_secs = all_end - all_start;
|
||||
|
||||
// read speed in GB/s
|
||||
double perf = (static_cast<double>(bufSize_ * numIter) *
|
||||
static_cast<double>(1e-09)) / elapsed_secs.count();
|
||||
|
||||
const char *strSrc = NULL;
|
||||
const char *strDst = NULL;
|
||||
if (hostMalloc[0])
|
||||
strSrc = "hHM";
|
||||
else if (hostRegister[0])
|
||||
strSrc = "hHR";
|
||||
else if (unpinnedMalloc[0])
|
||||
strSrc = "unp";
|
||||
else
|
||||
strSrc = "hM";
|
||||
|
||||
if (hostMalloc[1])
|
||||
strDst = "hHM";
|
||||
else if (hostRegister[1])
|
||||
strDst = "hHR";
|
||||
else if (unpinnedMalloc[1])
|
||||
strDst = "unp";
|
||||
else
|
||||
strDst = "hM";
|
||||
|
||||
// Double results when src and dst are both on device
|
||||
if ((!hostMalloc[0] && !hostRegister[0] && !unpinnedMalloc[0]) &&
|
||||
(!hostMalloc[1] && !hostRegister[1] && !unpinnedMalloc[1]))
|
||||
perf *= 2.0;
|
||||
// Double results when src and dst are both in sysmem
|
||||
if ((hostMalloc[0] || hostRegister[0] || unpinnedMalloc[0]) &&
|
||||
(hostMalloc[1] || hostRegister[1] || unpinnedMalloc[1]))
|
||||
perf *= 2.0;
|
||||
|
||||
INFO("hipPerfBufferCopyRectSpeed[" << test << "]\t( " << bufSize_ <<
|
||||
")\ts:" << strSrc << " d:" << strDst << "\ti:" << numIter <<
|
||||
"\t(GB/s) perf\t" << (float)perf);
|
||||
|
||||
// Free src
|
||||
if (hostMalloc[0]) {
|
||||
HIP_CHECK(hipHostFree(srcBuffer));
|
||||
} else if (hostRegister[0]) {
|
||||
HIP_CHECK(hipHostUnregister(srcBuffer));
|
||||
free(memptr[0]);
|
||||
} else if (unpinnedMalloc[0]) {
|
||||
free(memptr[0]);
|
||||
} else {
|
||||
HIP_CHECK(hipFree(srcBuffer));
|
||||
}
|
||||
|
||||
// Free dst
|
||||
if (hostMalloc[1]) {
|
||||
HIP_CHECK(hipHostFree(dstBuffer));
|
||||
} else if (hostRegister[1]) {
|
||||
HIP_CHECK(hipHostUnregister(dstBuffer));
|
||||
free(memptr[1]);
|
||||
} else if (unpinnedMalloc[1]) {
|
||||
free(memptr[1]);
|
||||
} else {
|
||||
HIP_CHECK(hipFree(dstBuffer));
|
||||
}
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipPerfBufferCopy status.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - perftests/memory/hipPerfBufferCopyRectSpeed.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.6
|
||||
*/
|
||||
|
||||
TEST_CASE("Perf_hipPerfBufferCopyRectSpeed_test") {
|
||||
int numDevices = 0;
|
||||
HIP_CHECK(hipGetDeviceCount(&numDevices));
|
||||
|
||||
if (numDevices <= 0) {
|
||||
SUCCEED("Skipped testcase hipPerfBufferCopyRectSpeed"
|
||||
"as there is no device to test.");
|
||||
} else {
|
||||
int deviceId = 0;
|
||||
HIP_CHECK(hipSetDevice(deviceId));
|
||||
hipDeviceProp_t props = {0};
|
||||
HIP_CHECK(hipGetDeviceProperties(&props, deviceId));
|
||||
|
||||
INFO("hipPerfBufferCopyRectSpeed - info: Set device to " << deviceId
|
||||
<< " : " << props.name << "Legend: unp - unpinned(malloc),"
|
||||
" hM - hipMalloc(device)\n hHR - hipHostRegister(pinned),"
|
||||
" hHM - hipHostMalloc(prePinned)\n");
|
||||
|
||||
REQUIRE(true == hipPerfBufferCopyRectSpeed_test(1));
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,258 @@
|
||||
/*
|
||||
Copyright (c) 2023 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.
|
||||
*/
|
||||
|
||||
/**
|
||||
* @addtogroup hipMemcpyAsync hipMemcpyAsync
|
||||
* @{
|
||||
* @ingroup perfMemoryTest
|
||||
* `hipMemcpyAsync(void* dst, const void* src, size_t count,
|
||||
* hipMemcpyKind kind, hipStream_t stream = 0)` -
|
||||
* Copies data between host and device.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
#define NUM_SIZES 9
|
||||
// 4KB, 8KB, 64KB, 256KB, 1 MB, 4MB, 16 MB, 16MB+10
|
||||
static const unsigned int Sizes[NUM_SIZES] =
|
||||
{4096, 8192, 65536, 262144, 524288, 1048576, 4194304, 16777216, 16777216+10};
|
||||
|
||||
static const unsigned int Iterations[2] = {1, 1000};
|
||||
|
||||
#define BUF_TYPES 4
|
||||
// 16 ways to combine 4 different buffer types
|
||||
#define NUM_SUBTESTS (BUF_TYPES*BUF_TYPES)
|
||||
|
||||
static void setData(void *ptr, unsigned int size, char value) {
|
||||
char *ptr2 = reinterpret_cast<char *>(ptr);
|
||||
for (unsigned int i = 0; i < size ; i++) {
|
||||
ptr2[i] = value;
|
||||
}
|
||||
}
|
||||
|
||||
static void checkData(void *ptr, unsigned int size, char value) {
|
||||
char *ptr2 = reinterpret_cast<char *>(ptr);
|
||||
for (unsigned int i = 0; i < size; i++) {
|
||||
if (ptr2[i] != value) {
|
||||
INFO("Validation failed at " << i << " Got " << ptr2[i] <<
|
||||
" Expected " << value);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
static bool hipPerfBufferCopySpeed_test(int p_tests) {
|
||||
unsigned int bufSize_;
|
||||
unsigned int numIter;
|
||||
bool hostMalloc[2] = {false};
|
||||
bool hostRegister[2] = {false};
|
||||
bool unpinnedMalloc[2] = {false};
|
||||
void *memptr[2] = {NULL};
|
||||
void *alignedmemptr[2] = {NULL};
|
||||
void *srcBuffer = NULL;
|
||||
void *dstBuffer = NULL;
|
||||
|
||||
int numTests = (p_tests == -1) ? (NUM_SIZES*NUM_SUBTESTS*2 - 1) : p_tests;
|
||||
int test = (p_tests == -1) ? 0 : p_tests;
|
||||
|
||||
for ( ; test <= numTests; test++ ) {
|
||||
unsigned int srcTest = (test / NUM_SIZES) % BUF_TYPES;
|
||||
unsigned int dstTest = (test / (NUM_SIZES*BUF_TYPES)) % BUF_TYPES;
|
||||
bufSize_ = Sizes[test % NUM_SIZES];
|
||||
hostMalloc[0] = hostMalloc[1] = false;
|
||||
hostRegister[0] = hostRegister[1] = false;
|
||||
unpinnedMalloc[0] = unpinnedMalloc[1] = false;
|
||||
srcBuffer = dstBuffer = 0;
|
||||
memptr[0] = memptr[1] = NULL;
|
||||
alignedmemptr[0] = alignedmemptr[1] = NULL;
|
||||
|
||||
if (srcTest == 3) {
|
||||
hostRegister[0] = true;
|
||||
} else if (srcTest == 2) {
|
||||
hostMalloc[0] = true;
|
||||
} else if (srcTest == 1) {
|
||||
unpinnedMalloc[0] = true;
|
||||
}
|
||||
|
||||
if (dstTest == 1) {
|
||||
unpinnedMalloc[1] = true;
|
||||
} else if (dstTest == 2) {
|
||||
hostMalloc[1] = true;
|
||||
} else if (dstTest == 3) {
|
||||
hostRegister[1] = true;
|
||||
}
|
||||
|
||||
numIter = Iterations[test / (NUM_SIZES * NUM_SUBTESTS)];
|
||||
|
||||
if (hostMalloc[0]) {
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&srcBuffer),
|
||||
bufSize_, 0));
|
||||
setData(srcBuffer, bufSize_, 0xd0);
|
||||
} else if (hostRegister[0]) {
|
||||
memptr[0] = malloc(bufSize_ + 4096);
|
||||
alignedmemptr[0] = reinterpret_cast<void*>(memptr[0]);
|
||||
srcBuffer = alignedmemptr[0];
|
||||
setData(srcBuffer, bufSize_, 0xd0);
|
||||
HIP_CHECK(hipHostRegister(srcBuffer, bufSize_, 0));
|
||||
} else if (unpinnedMalloc[0]) {
|
||||
memptr[0] = malloc(bufSize_ + 4096);
|
||||
alignedmemptr[0] = reinterpret_cast<void*>(memptr[0]);
|
||||
srcBuffer = alignedmemptr[0];
|
||||
setData(srcBuffer, bufSize_, 0xd0);
|
||||
} else {
|
||||
HIP_CHECK(hipMalloc(&srcBuffer, bufSize_));
|
||||
HIP_CHECK(hipMemset(srcBuffer, 0xd0, bufSize_));
|
||||
}
|
||||
|
||||
if (hostMalloc[1]) {
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&dstBuffer),
|
||||
bufSize_, 0));
|
||||
} else if (hostRegister[1]) {
|
||||
memptr[1] = malloc(bufSize_ + 4096);
|
||||
alignedmemptr[1] = reinterpret_cast<void*>(memptr[1]);
|
||||
dstBuffer = alignedmemptr[1];
|
||||
HIP_CHECK(hipHostRegister(dstBuffer, bufSize_, 0));
|
||||
} else if (unpinnedMalloc[1]) {
|
||||
memptr[1] = malloc(bufSize_ + 4096);
|
||||
alignedmemptr[1] = reinterpret_cast<void*>(memptr[1]);
|
||||
dstBuffer = alignedmemptr[1];
|
||||
} else {
|
||||
HIP_CHECK(hipMalloc(&dstBuffer, bufSize_));
|
||||
}
|
||||
|
||||
// warm up
|
||||
HIP_CHECK(hipMemcpy(dstBuffer, srcBuffer, bufSize_, hipMemcpyDefault));
|
||||
|
||||
// measure performance based on host time
|
||||
auto all_start = std::chrono::steady_clock::now();
|
||||
|
||||
for (unsigned int i = 0; i < numIter; i++) {
|
||||
HIP_CHECK(hipMemcpyAsync(dstBuffer, srcBuffer, bufSize_,
|
||||
hipMemcpyDefault, NULL));
|
||||
}
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
auto all_end = std::chrono::steady_clock::now();
|
||||
std::chrono::duration<double> elapsed_secs = all_end - all_start;
|
||||
|
||||
// read speed in GB/s
|
||||
double perf = (static_cast<double>(bufSize_ * numIter) *
|
||||
static_cast<double>(1e-09)) / elapsed_secs.count();
|
||||
|
||||
const char *strSrc = NULL;
|
||||
const char *strDst = NULL;
|
||||
if (hostMalloc[0])
|
||||
strSrc = "hHM";
|
||||
else if (hostRegister[0])
|
||||
strSrc = "hHR";
|
||||
else if (unpinnedMalloc[0])
|
||||
strSrc = "unp";
|
||||
else
|
||||
strSrc = "hM";
|
||||
|
||||
if (hostMalloc[1])
|
||||
strDst = "hHM";
|
||||
else if (hostRegister[1])
|
||||
strDst = "hHR";
|
||||
else if (unpinnedMalloc[1])
|
||||
strDst = "unp";
|
||||
else
|
||||
strDst = "hM";
|
||||
|
||||
// Double results when src and dst are both on device
|
||||
if ((!hostMalloc[0] && !hostRegister[0] && !unpinnedMalloc[0]) &&
|
||||
(!hostMalloc[1] && !hostRegister[1] && !unpinnedMalloc[1]))
|
||||
perf *= 2.0;
|
||||
// Double results when src and dst are both in sysmem
|
||||
if ((hostMalloc[0] || hostRegister[0] || unpinnedMalloc[0]) &&
|
||||
(hostMalloc[1] || hostRegister[1] || unpinnedMalloc[1]))
|
||||
perf *= 2.0;
|
||||
|
||||
INFO("HIPPerfBufferCopySpeed[" << test << "]\t( " << bufSize_ <<
|
||||
")\ts:" << strSrc << " d:" << strDst << "\ti:" << numIter <<
|
||||
"\t(GB/s) perf\t" << (float)perf);
|
||||
|
||||
// Verification
|
||||
void* temp = malloc(bufSize_ + 4096);
|
||||
void* chkBuf = reinterpret_cast<void*>(temp);
|
||||
HIP_CHECK(hipMemcpy(chkBuf, dstBuffer, bufSize_, hipMemcpyDefault));
|
||||
checkData(chkBuf, bufSize_, 0xd0);
|
||||
free(temp);
|
||||
|
||||
// Free src
|
||||
if (hostMalloc[0]) {
|
||||
HIP_CHECK(hipHostFree(srcBuffer));
|
||||
} else if (hostRegister[0]) {
|
||||
HIP_CHECK(hipHostUnregister(srcBuffer));
|
||||
free(memptr[0]);
|
||||
} else if (unpinnedMalloc[0]) {
|
||||
free(memptr[0]);
|
||||
} else {
|
||||
HIP_CHECK(hipFree(srcBuffer));
|
||||
}
|
||||
|
||||
// Free dst
|
||||
if (hostMalloc[1]) {
|
||||
HIP_CHECK(hipHostFree(dstBuffer));
|
||||
} else if (hostRegister[1]) {
|
||||
HIP_CHECK(hipHostUnregister(dstBuffer));
|
||||
free(memptr[1]);
|
||||
} else if (unpinnedMalloc[1]) {
|
||||
free(memptr[1]);
|
||||
} else {
|
||||
HIP_CHECK(hipFree(dstBuffer));
|
||||
}
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipPerfBufferCopySpeed status.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - perftests/memory/hipPerfBufferCopySpeed.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.6
|
||||
*/
|
||||
|
||||
TEST_CASE("Perf_hipPerfBufferCopySpeed_test") {
|
||||
int numDevices = 0;
|
||||
HIP_CHECK(hipGetDeviceCount(&numDevices));
|
||||
|
||||
if (numDevices <= 0) {
|
||||
SUCCEED("Skipped testcase hipPerfBufferCopySpeed as"
|
||||
"there is no device to test.");
|
||||
} else {
|
||||
int deviceId = 0;
|
||||
HIP_CHECK(hipSetDevice(deviceId));
|
||||
hipDeviceProp_t props = {0};
|
||||
HIP_CHECK(hipGetDeviceProperties(&props, deviceId));
|
||||
|
||||
INFO("hipPerfBufferCopySpeed - info: Set device to " << deviceId
|
||||
<< " : " << props.name << "Legend: unp - unpinned(malloc),"
|
||||
" hM - hipMalloc(device)\n hHR - hipHostRegister(pinned),"
|
||||
" hHM - hipHostMalloc(prePinned)\n");
|
||||
|
||||
REQUIRE(true == hipPerfBufferCopySpeed_test(1));
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,154 @@
|
||||
/*
|
||||
Copyright (c) 2023 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.
|
||||
*/
|
||||
|
||||
/**
|
||||
* @addtogroup hipMemcpyKernel hipMemcpyKernel
|
||||
* @{
|
||||
* @ingroup perfMemoryTest
|
||||
* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` -
|
||||
* Copies data between host and device.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
#define ARRAY_SIZE 16
|
||||
|
||||
typedef struct d_uint16 {
|
||||
uint data[ARRAY_SIZE];
|
||||
} d_uint16;
|
||||
|
||||
__global__ static void read_kernel(d_uint16 *src, ulong N, uint *dst) {
|
||||
size_t idx = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
size_t stride = blockDim.x * gridDim.x;
|
||||
|
||||
uint tmp = 0;
|
||||
for (size_t i = idx; i < N; i += stride) {
|
||||
for (size_t j = 0; j < ARRAY_SIZE; j++) {
|
||||
tmp += src[i].data[j];
|
||||
}
|
||||
}
|
||||
|
||||
atomicAdd(dst, tmp);
|
||||
}
|
||||
|
||||
static bool hipPerfDevMemReadSpeed_test() {
|
||||
d_uint16 *dSrc, *hSrc;
|
||||
uint *dDst, *hDst;
|
||||
hipStream_t stream;
|
||||
ulong N = 4 * 1024 * 1024;
|
||||
uint nBytes = N * sizeof(d_uint16);
|
||||
|
||||
int deviceId = 0;
|
||||
HIP_CHECK(hipSetDevice(deviceId));
|
||||
hipDeviceProp_t props = {0};
|
||||
HIP_CHECK(hipGetDeviceProperties(&props, deviceId));
|
||||
|
||||
INFO("info: running on bus " << "0x" << props.pciBusID << " " <<
|
||||
props.name << " with " << props.multiProcessorCount << " CUs \n");
|
||||
|
||||
const unsigned threadsPerBlock = 64;
|
||||
const unsigned blocks = props.multiProcessorCount * 4;
|
||||
|
||||
uint inputData = 0x1;
|
||||
int nIter = 1000;
|
||||
|
||||
hSrc = new d_uint16[nBytes];
|
||||
REQUIRE(hSrc != nullptr);
|
||||
hDst = new uint;
|
||||
REQUIRE(hDst != nullptr);
|
||||
hDst[0] = 0;
|
||||
|
||||
for (size_t i = 0; i < N; i++) {
|
||||
for (int j = 0; j < ARRAY_SIZE; j++) {
|
||||
hSrc[i].data[j] = inputData;
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipMalloc(&dSrc, nBytes));
|
||||
HIP_CHECK(hipMalloc(&dDst, sizeof(uint)));
|
||||
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIP_CHECK(hipMemcpy(dSrc, hSrc, nBytes, hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipMemcpy(dDst, hDst, sizeof(uint), hipMemcpyHostToDevice));
|
||||
|
||||
hipLaunchKernelGGL(read_kernel, dim3(blocks), dim3(threadsPerBlock),
|
||||
0, stream, dSrc, N, dDst);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
HIP_CHECK(hipMemcpy(hDst, dDst, sizeof(uint), hipMemcpyDeviceToHost));
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
if (hDst[0] != (nBytes / sizeof(uint))) {
|
||||
INFO("hipPerfDevMemReadSpeed - Data validation failed for warm up run!" <<
|
||||
" expected " << nBytes / sizeof(uint) << " got " << hDst[0]);
|
||||
return false;
|
||||
}
|
||||
|
||||
// measure performance based on host time
|
||||
auto all_start = std::chrono::steady_clock::now();
|
||||
|
||||
for (int i = 0; i < nIter; i++) {
|
||||
hipLaunchKernelGGL(read_kernel, dim3(blocks), dim3(threadsPerBlock),
|
||||
0, stream, dSrc, N, dDst);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
}
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
auto all_end = std::chrono::steady_clock::now();
|
||||
std::chrono::duration<double> all_kernel_time = all_end - all_start;
|
||||
|
||||
// read speed in GB/s
|
||||
double perf = (static_cast<double>(nBytes * nIter * (1e-09))) /
|
||||
all_kernel_time.count();
|
||||
|
||||
INFO("hipPerfDevMemReadSpeed - info: average read speed of " <<
|
||||
perf << " GB/s " << "achieved for memory size of " <<
|
||||
nBytes / (1024 * 1024) << " MB");
|
||||
|
||||
delete [] hSrc;
|
||||
delete hDst;
|
||||
HIP_CHECK(hipFree(dSrc));
|
||||
HIP_CHECK(hipFree(dDst));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
return true;
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipPerfDevMemReadSpeed status.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - perftests/memory/hipPerfDevMemReadSpeed.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.6
|
||||
*/
|
||||
|
||||
TEST_CASE("Perf_hipPerfDevMemReadSpeed_test") {
|
||||
int numDevices = 0;
|
||||
HIP_CHECK(hipGetDeviceCount(&numDevices));
|
||||
|
||||
if (numDevices <= 0) {
|
||||
SUCCEED("Skipped testcase hipPerfDevMemReadSpeed as"
|
||||
"there is no device to test.");
|
||||
} else {
|
||||
REQUIRE(true == hipPerfDevMemReadSpeed_test());
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,147 @@
|
||||
/*
|
||||
Copyright (c) 2023 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.
|
||||
*/
|
||||
|
||||
/**
|
||||
* @addtogroup hipMemcpyKernel hipMemcpyKernel
|
||||
* @{
|
||||
* @ingroup perfMemoryTest
|
||||
* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` -
|
||||
* Copies data between host and device.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
#define ARRAY_SIZE 16
|
||||
|
||||
typedef struct d_uint16 {
|
||||
uint data[ARRAY_SIZE];
|
||||
} d_uint16;
|
||||
|
||||
__global__ void write_kernel(d_uint16 *dst, ulong N, d_uint16 pval) {
|
||||
size_t idx = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
size_t stride = blockDim.x * gridDim.x;
|
||||
for (size_t i = idx; i < N; i += stride) {
|
||||
dst[i] = pval;
|
||||
}
|
||||
}
|
||||
|
||||
static bool hipPerfDevMemWriteSpeed_test() {
|
||||
d_uint16 *dDst, *hDst;
|
||||
ulong N = 4 * 1024 * 1024;
|
||||
uint nBytes = N * sizeof(d_uint16);
|
||||
|
||||
uint inputData = 0xabababab;
|
||||
int nIter = 1000;
|
||||
d_uint16 pval;
|
||||
|
||||
int deviceId = 0;
|
||||
HIP_CHECK(hipSetDevice(deviceId));
|
||||
hipDeviceProp_t props = {0};
|
||||
HIP_CHECK(hipGetDeviceProperties(&props, deviceId));
|
||||
|
||||
INFO("info: running on bus " << "0x" << props.pciBusID << " " <<
|
||||
props.name << " with " << props.multiProcessorCount << " CUs \n");
|
||||
|
||||
const unsigned threadsPerBlock = 64;
|
||||
const unsigned blocks = props.multiProcessorCount * 4;
|
||||
|
||||
for (int i = 0; i < ARRAY_SIZE; i++) {
|
||||
pval.data[i] = inputData;
|
||||
}
|
||||
|
||||
hDst = new d_uint16[nBytes];
|
||||
REQUIRE(hDst != nullptr);
|
||||
|
||||
for (size_t i = 0; i < N; i++) {
|
||||
for (size_t j = 0; j < ARRAY_SIZE; j++) {
|
||||
hDst[i].data[j] = 0;
|
||||
}
|
||||
}
|
||||
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIP_CHECK(hipMalloc(&dDst, nBytes));
|
||||
hipLaunchKernelGGL(write_kernel, dim3(blocks), dim3(threadsPerBlock),
|
||||
0, stream, dDst, N, pval);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
HIP_CHECK(hipMemcpy(hDst, dDst, nBytes , hipMemcpyDeviceToHost));
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
for (uint i = 0; i < N; i++) {
|
||||
for (uint j = 0; j < ARRAY_SIZE; j++) {
|
||||
if (hDst[i].data[j] != inputData) {
|
||||
INFO("hipPerfDevMemWriteSpeed - Data validation failed for warm up run!"
|
||||
<< "at index i: " << i << " element j: " << j <<
|
||||
"expected " << inputData << " but got " << hDst[i].data[j]);
|
||||
return false;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// measure performance based on host time
|
||||
auto all_start = std::chrono::steady_clock::now();
|
||||
|
||||
for (int i = 0; i < nIter; i++) {
|
||||
hipLaunchKernelGGL(write_kernel, dim3(blocks), dim3(threadsPerBlock),
|
||||
0, stream, dDst, N, pval);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
}
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
auto all_end = std::chrono::steady_clock::now();
|
||||
std::chrono::duration<double> all_kernel_time = all_end - all_start;
|
||||
|
||||
// read speed in GB/s
|
||||
double perf = (static_cast<double>(nBytes * nIter * (1e-09))) /
|
||||
all_kernel_time.count();
|
||||
|
||||
INFO("hipPerfDevMemReadSpeed - info: average write speed of " <<
|
||||
perf << " GB/s " << "achieved for memory size of " <<
|
||||
nBytes / (1024 * 1024) << " MB");
|
||||
|
||||
delete [] hDst;
|
||||
HIP_CHECK(hipFree(dDst));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
return true;
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipPerfDevMemWriteSpeed status.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - perftests/memory/hipPerfDevMemWriteSpeed.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.6
|
||||
*/
|
||||
|
||||
TEST_CASE("Perf_hipPerfDevMemWriteSpeed_test") {
|
||||
int numDevices = 0;
|
||||
HIP_CHECK(hipGetDeviceCount(&numDevices));
|
||||
|
||||
if (numDevices <= 0) {
|
||||
SUCCEED("Skipped testcase hipPerfDevMemWriteSpeed as"
|
||||
"there is no device to test.");
|
||||
} else {
|
||||
REQUIRE(true == hipPerfDevMemWriteSpeed_test());
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,191 @@
|
||||
/*
|
||||
Copyright (c) 2023 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.
|
||||
*/
|
||||
|
||||
/**
|
||||
* @addtogroup hipMemcpyKernel hipMemcpyKernel
|
||||
* @{
|
||||
* @ingroup perfMemoryTest
|
||||
* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` -
|
||||
* Copies data between host and device.
|
||||
*/
|
||||
|
||||
#include <numaif.h>
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
// To run it correctly, we must not export HIP_VISIBLE_DEVICES.
|
||||
// And we must explicitly link libnuma because of numa api move_pages().
|
||||
#define NUM_PAGES 4
|
||||
char *h = nullptr;
|
||||
char *d_h = nullptr;
|
||||
char *m = nullptr;
|
||||
char *d_m = nullptr;
|
||||
int page_size = 1024;
|
||||
|
||||
const int mode[] = { MPOL_DEFAULT, MPOL_BIND, MPOL_PREFERRED, MPOL_INTERLEAVE };
|
||||
const char* modeStr[] = { "MPOL_DEFAULT", "MPOL_BIND",
|
||||
"MPOL_PREFERRED", "MPOL_INTERLEAVE" };
|
||||
|
||||
std::string exeCommand(const char* cmd) {
|
||||
std::array<char, 128> buff;
|
||||
std::string result;
|
||||
std::unique_ptr<FILE, decltype(&pclose)> pipe(popen(cmd, "r"), pclose);
|
||||
if (!pipe) {
|
||||
return result;
|
||||
}
|
||||
while (fgets(buff.data(), buff.size(), pipe.get()) != nullptr) {
|
||||
result += buff.data();
|
||||
}
|
||||
return result;
|
||||
}
|
||||
|
||||
int getCpuAgentCount() {
|
||||
const char* cmd =
|
||||
"cat /proc/cpuinfo | grep \"physical id\" | sort | uniq | wc -l";
|
||||
int cpuAgentCount = std::atoi(exeCommand(cmd).c_str());
|
||||
return cpuAgentCount;
|
||||
}
|
||||
|
||||
bool test(int cpuId, int gpuId, int numaMode, unsigned int hostMallocflags) {
|
||||
void *pages[NUM_PAGES];
|
||||
int status[NUM_PAGES];
|
||||
int nodes[NUM_PAGES];
|
||||
int ret_code;
|
||||
|
||||
INFO("set cpu " << cpuId << ", gpu " << gpuId << ", numaMode "
|
||||
<< numaMode << ", hostMallocflags " << hostMallocflags << "\n");
|
||||
|
||||
if (cpuId >= 0) {
|
||||
unsigned long nodeMask = 1 << cpuId; //NOLINT
|
||||
unsigned long maxNode = sizeof(nodeMask) * 8; //NOLINT
|
||||
if (set_mempolicy(numaMode, numaMode == MPOL_DEFAULT ? NULL : &nodeMask,
|
||||
numaMode == MPOL_DEFAULT ? 0 : maxNode) == -1) {
|
||||
WARN("set_mempolicy() failed with err " << errno << "\n");
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
if (gpuId >= 0) {
|
||||
HIP_CHECK(hipSetDevice(gpuId));
|
||||
}
|
||||
|
||||
posix_memalign(reinterpret_cast<void**>(&m), page_size, page_size*NUM_PAGES);
|
||||
HIP_CHECK(hipHostRegister(m, page_size * NUM_PAGES, hipHostRegisterMapped));
|
||||
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&d_m), m, 0));
|
||||
|
||||
status[0] = -1;
|
||||
pages[0] = m;
|
||||
for (int i = 1; i < NUM_PAGES; i++) {
|
||||
pages[i] = reinterpret_cast<char*>(pages[0]) + page_size;
|
||||
}
|
||||
|
||||
ret_code = move_pages(0, NUM_PAGES, pages, NULL, status, 0);
|
||||
INFO("Memory (malloc) ret " << ret_code << " at " << m <<
|
||||
" (dev " << d_m << "%p) is at node: ");
|
||||
for (int i = 0; i < NUM_PAGES; i++) {
|
||||
INFO(status[i]); // Don't verify as it's out of our control
|
||||
}
|
||||
INFO("\n");
|
||||
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&h),
|
||||
page_size*NUM_PAGES, hostMallocflags));
|
||||
pages[0] = h;
|
||||
for (int i = 1; i < NUM_PAGES; i++) {
|
||||
pages[i] = reinterpret_cast<char*>(pages[0]) + page_size;
|
||||
}
|
||||
ret_code = move_pages(0, NUM_PAGES, pages, NULL, status, 0);
|
||||
d_h = nullptr;
|
||||
if (hostMallocflags & hipHostMallocMapped) {
|
||||
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&d_h), h, 0));
|
||||
INFO("Memory (hipHostMalloc) ret " << ret_code << " at " << h
|
||||
<< " (dev " << d_h << ") is at node: ");
|
||||
} else {
|
||||
INFO("Memory (hipHostMalloc) ret " << ret_code << " at "
|
||||
<< h << " is at node: ");
|
||||
}
|
||||
for (int i = 0; i < NUM_PAGES; i++) {
|
||||
INFO(status[i]); // Always print it even if it's wrong. Verify later
|
||||
}
|
||||
INFO("\n");
|
||||
|
||||
HIP_CHECK(hipHostFree(reinterpret_cast<void*>(h)));
|
||||
HIP_CHECK(hipHostUnregister(m));
|
||||
free(m);
|
||||
|
||||
if (cpuId >= 0 && (numaMode == MPOL_BIND || numaMode == MPOL_PREFERRED)) {
|
||||
for (int i = 0; i < NUM_PAGES; i++) {
|
||||
if (status[i] != cpuId) { // Now verify
|
||||
WARN("Failed at " << i << " status[i] = " << status[i]
|
||||
<< " cpuId " << cpuId << "\n");
|
||||
return false;
|
||||
}
|
||||
}
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
bool runTest(const int &cpuCount, const int &gpuCount,
|
||||
unsigned int hostMallocflags, const std::string &str) {
|
||||
INFO("Test- " << str.c_str() << "\n");
|
||||
|
||||
for (int m = 0; m < sizeof(mode) / sizeof(mode[0]); m++) {
|
||||
INFO("Testing " << modeStr[m] << "\n");
|
||||
|
||||
for (int i = 0; i < cpuCount; i++) {
|
||||
for (int j = 0; j < gpuCount; j++) {
|
||||
if (!test(i, j, mode[m], hostMallocflags)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipPerfHostNumaAlloc status.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - perftests/memory/hipPerfHostNumaAlloc.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.6
|
||||
*/
|
||||
|
||||
TEST_CASE("Perf_hipPerfHostNumaAlloc_test") {
|
||||
int gpuCount = 0;
|
||||
HIP_CHECK(hipGetDeviceCount(&gpuCount));
|
||||
int cpuCount = getCpuAgentCount();
|
||||
INFO("Cpu count " << cpuCount << ", Gpu count " << gpuCount << "\n");
|
||||
|
||||
if (cpuCount < 0 || gpuCount < 0) {
|
||||
SUCCEED("Skipped testcase hipPerfHostNumaAlloc as "
|
||||
"there is no device to test.\n");
|
||||
return;
|
||||
}
|
||||
|
||||
REQUIRE(true == runTest(cpuCount, gpuCount,
|
||||
hipHostMallocDefault | hipHostMallocNumaUser,
|
||||
"Testing hipHostMallocDefault | hipHostMallocNumaUser......"));
|
||||
|
||||
REQUIRE(true == runTest(cpuCount, gpuCount,
|
||||
hipHostMallocMapped | hipHostMallocNumaUser,
|
||||
"Testing hipHostMallocMapped | hipHostMallocNumaUser......."));
|
||||
}
|
||||
@@ -0,0 +1,545 @@
|
||||
/*
|
||||
Copyright (c) 2023 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.
|
||||
*/
|
||||
|
||||
/**
|
||||
* @addtogroup hipMemcpyKernel hipMemcpyKernel
|
||||
* @{
|
||||
* @ingroup perfMemoryTest
|
||||
* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` -
|
||||
* Copies data between host and device.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
#define SIMPLY_ASSIGN 0
|
||||
#define USE_HIPTEST_SETNUMBLOCKS 0
|
||||
|
||||
template<class T>
|
||||
__global__ void vec_fill(T *x, T coef, int N) {
|
||||
const int istart = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
const int ishift = blockDim.x * gridDim.x;
|
||||
for (int i = istart; i < N; i += ishift) {
|
||||
#if SIMPLY_ASSIGN
|
||||
x[i] = coef;
|
||||
#else
|
||||
x[i] = coef * i;
|
||||
#endif
|
||||
}
|
||||
}
|
||||
|
||||
__device__ void print_log(int i, double value, double expected) {
|
||||
printf("failed at %d: val=%g, expected=%g\n", i, value, expected);
|
||||
}
|
||||
|
||||
__device__ void print_log(int i, int value, int expected) {
|
||||
printf("failed at %d: val=%d, expected=%d\n", i, value, expected);
|
||||
}
|
||||
|
||||
template<class T>
|
||||
__global__ void vec_verify(T *x, T coef, int N) {
|
||||
const int istart = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
const int ishift = blockDim.x * gridDim.x;
|
||||
for (int i = istart; i < N; i += ishift) {
|
||||
#if SIMPLY_ASSIGN
|
||||
if (x[i] != coef) {
|
||||
print_log(i, x[i], coef);
|
||||
}
|
||||
#else
|
||||
if (x[i] != coef * i) {
|
||||
print_log(i, x[i], coef * i);
|
||||
}
|
||||
#endif
|
||||
}
|
||||
}
|
||||
|
||||
template<class T>
|
||||
__global__ void daxpy(T *__restrict__ x, T *__restrict__ y,
|
||||
const T coef, int Niter, int N) {
|
||||
const int istart = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
const int ishift = blockDim.x * gridDim.x;
|
||||
for (int iter = 0; iter < Niter; ++iter) {
|
||||
T iv = coef * iter;
|
||||
for (int i = istart; i < N; i += ishift)
|
||||
y[i] = iv * x[i] + y[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<class T>
|
||||
class hipPerfMemFill {
|
||||
private:
|
||||
static constexpr int NUM_START = 27;
|
||||
static constexpr int NUM_SIZE = 5;
|
||||
static constexpr int NUM_ITER = 10;
|
||||
static constexpr double NUM_1GB = 1024.0 * 1024.0 * 1024.0;
|
||||
size_t totalSizes_[NUM_SIZE];
|
||||
hipDeviceProp_t props_;
|
||||
const T coef_ = getCoefficient(3.14159);
|
||||
const unsigned int threadsPerBlock_ = 64;
|
||||
unsigned int blocksPerCU_;
|
||||
|
||||
public:
|
||||
hipPerfMemFill() {
|
||||
for (int i = 0; i < NUM_SIZE; i++) {
|
||||
// 128M, 256M, 512M, 1024M, 2048M
|
||||
totalSizes_[i] = 1ull << (i + NUM_START);
|
||||
}
|
||||
}
|
||||
|
||||
~hipPerfMemFill() { }
|
||||
|
||||
bool supportLargeBar() {
|
||||
return props_.isLargeBar != 0;
|
||||
}
|
||||
|
||||
bool supportManagedMemory() {
|
||||
return props_.managedMemory != 0;
|
||||
}
|
||||
|
||||
const T getCoefficient(double val) {
|
||||
return static_cast<T>(val);
|
||||
}
|
||||
|
||||
void setHostBuffer(T *A, T val, size_t size) {
|
||||
size_t len = size / sizeof(T);
|
||||
for (int i = 0; i < len; i++) {
|
||||
A[i] = val;
|
||||
}
|
||||
}
|
||||
|
||||
bool open(int deviceId) {
|
||||
int nGpu = 0;
|
||||
HIP_CHECK(hipGetDeviceCount(&nGpu));
|
||||
if (nGpu < 1) {
|
||||
printf("No GPU!");
|
||||
return false;
|
||||
} else if (deviceId >= nGpu) {
|
||||
printf("Info: wrong GPU Id %d\n", deviceId);
|
||||
return false;
|
||||
}
|
||||
|
||||
HIP_CHECK(hipSetDevice(deviceId));
|
||||
memset(&props_, 0, sizeof(props_));
|
||||
HIP_CHECK(hipGetDeviceProperties(&props_, deviceId));
|
||||
blocksPerCU_ = props_.multiProcessorCount * 4;
|
||||
|
||||
std::cout << "Info: running on device: id: " << deviceId << ", bus: 0x"
|
||||
<< props_.pciBusID << " " << props_.name << " with "
|
||||
<< props_.multiProcessorCount << " CUs, large bar: "
|
||||
<< supportLargeBar() << ", managed memory: " << supportManagedMemory()
|
||||
<< ", DeviceMallocFinegrained: " << supportDeviceMallocFinegrained()
|
||||
<< std::endl;
|
||||
return true;
|
||||
}
|
||||
|
||||
void log_host(const char* title, double GBytes, double sec) {
|
||||
std::cout << title << " [" << std::setw(7) << GBytes << " GB]: cost "
|
||||
<< std::setw(10) << sec << " s in bandwidth " << std::setw(10)
|
||||
<< GBytes / sec << " [GB/s]" << std::endl;
|
||||
}
|
||||
|
||||
void log_kernel(const char* title, double GBytes, double sec,
|
||||
double sec_hv, double sec_kv) {
|
||||
std::cout << title << " [" << std::setw(7) << GBytes << " GB]: cost "
|
||||
<< std::setw(10) << sec << " s in bandwidth " << std::setw(10)
|
||||
<< GBytes / sec << " [GB/s]" << ", hostVerify cost "
|
||||
<< std::setw(10) << sec_hv << " s in bandwidth " << std::setw(10)
|
||||
<< GBytes / sec_hv << " [GB/s]" << ", kernelVerify cost "
|
||||
<< std::setw(10) << sec_kv << " s in bandwidth " << std::setw(10)
|
||||
<< GBytes / sec_kv << " [GB/s]" << std::endl;
|
||||
}
|
||||
|
||||
void hostFill(size_t size, T *data, T coef, double *sec) {
|
||||
size_t num = size / sizeof(T); // Size of elements
|
||||
auto start = std::chrono::steady_clock::now();
|
||||
for (int i = 0; i < num; ++i) {
|
||||
#if SIMPLY_ASSIGN
|
||||
data[i] = coef;
|
||||
#else
|
||||
data[i] = coef * i;
|
||||
#endif
|
||||
}
|
||||
auto end = std::chrono::steady_clock::now();
|
||||
std::chrono::duration<double> diff = end - start; // in second
|
||||
*sec = diff.count();
|
||||
}
|
||||
|
||||
void kernelFill(size_t size, T *data, T coef, double *sec) {
|
||||
size_t num = size / sizeof(T); // Size of elements
|
||||
unsigned blocks = setNumBlocks(num);
|
||||
|
||||
// kernel will be loaded first time
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_fill<T>), dim3(blocks),
|
||||
dim3(threadsPerBlock_), 0, 0, data, 0, num);
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
auto start = std::chrono::steady_clock::now();
|
||||
|
||||
for (int iter = 0; iter < NUM_ITER; ++iter) {
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_fill<T>), dim3(blocks),
|
||||
dim3(threadsPerBlock_), 0, 0, data, coef, num);
|
||||
}
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
auto end = std::chrono::steady_clock::now();
|
||||
std::chrono::duration<double> diff = end - start; // in second
|
||||
*sec = diff.count() / NUM_ITER; // in second
|
||||
}
|
||||
|
||||
void hostVerify(size_t size, T *data, T coef, double *sec) {
|
||||
size_t num = size / sizeof(T); // Size of elements
|
||||
auto start = std::chrono::steady_clock::now();
|
||||
for (int i = 0; i < num; ++i) {
|
||||
#if SIMPLY_ASSIGN
|
||||
if (data[i] != coef) {
|
||||
std::cout << "hostVerify failed: i=" << i << ", data[i]=" << data[i]
|
||||
<< ", expected=" << coef << std::endl;
|
||||
REQUIRE(false);
|
||||
}
|
||||
#else
|
||||
if (data[i] != coef * i) {
|
||||
std::cout << "hostVerify failed: i=" << i << ", data[i]=" << data[i]
|
||||
<< ", expected=" << coef * i << std::endl;
|
||||
REQUIRE(false);
|
||||
}
|
||||
#endif
|
||||
}
|
||||
auto end = std::chrono::steady_clock::now();
|
||||
std::chrono::duration<double> diff = end - start; // in second
|
||||
*sec = diff.count();
|
||||
}
|
||||
|
||||
void kernelVerify(size_t size, T *data, T coef, double *sec) {
|
||||
size_t num = size / sizeof(T); // Size of elements
|
||||
unsigned blocks = setNumBlocks(num);
|
||||
|
||||
// kernel will be loaded first time
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_verify<T>), dim3(blocks),
|
||||
dim3(threadsPerBlock_), 0, 0, data, coef, num);
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
// Now all data verified. The following is to test bandwidth.
|
||||
auto start = std::chrono::steady_clock::now();
|
||||
|
||||
for (int iter = 0; iter < NUM_ITER; ++iter) {
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_verify<T>), dim3(blocks),
|
||||
dim3(threadsPerBlock_), 0, 0, data, coef, num);
|
||||
}
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
auto end = std::chrono::steady_clock::now();
|
||||
std::chrono::duration<double> diff = end - start; // in second
|
||||
*sec = diff.count() / NUM_ITER; // in second
|
||||
}
|
||||
|
||||
bool testLargeBarDeviceMemoryHostFill(size_t size) {
|
||||
if (!supportLargeBar()) {
|
||||
return false;
|
||||
}
|
||||
|
||||
double GBytes = static_cast<double>(size) / NUM_1GB;
|
||||
|
||||
T *A;
|
||||
HIP_CHECK(hipMalloc(&A, size));
|
||||
double sec = 0;
|
||||
hostFill(size, A, coef_, &sec); // Cpu can access device mem in LB
|
||||
HIP_CHECK(hipFree(A));
|
||||
|
||||
log_host("Largebar: host fill", GBytes, sec);
|
||||
return true;
|
||||
}
|
||||
|
||||
bool testLargeBar() {
|
||||
if (!supportLargeBar()) {
|
||||
return false;
|
||||
}
|
||||
|
||||
std::cout << "Test large bar device memory host filling" << std::endl;
|
||||
for (int i = 0; i < NUM_SIZE; i++) {
|
||||
if (!testLargeBarDeviceMemoryHostFill(totalSizes_[i])) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool testManagedMemoryHostFill(size_t size) {
|
||||
if (!supportManagedMemory()) {
|
||||
return false;
|
||||
}
|
||||
double GBytes = static_cast<double>(size) / NUM_1GB;
|
||||
|
||||
T *A;
|
||||
HIP_CHECK(hipMallocManaged(&A, size));
|
||||
double sec = 0;
|
||||
hostFill(size, A, coef_, &sec); // Cpu can access HMM mem
|
||||
HIP_CHECK(hipFree(A));
|
||||
|
||||
log_host("Managed: host fill", GBytes, sec);
|
||||
return true;
|
||||
}
|
||||
|
||||
bool testManagedMemoryKernelFill(size_t size) {
|
||||
if (!supportManagedMemory()) {
|
||||
return false;
|
||||
}
|
||||
double GBytes = static_cast<double>(size) / NUM_1GB;
|
||||
|
||||
T *A;
|
||||
HIP_CHECK(hipMallocManaged(&A, size));
|
||||
|
||||
double sec = 0, sec_hv = 0, sec_kv = 0;
|
||||
kernelFill(size, A, coef_, &sec);
|
||||
// Managed memory can be verified by host
|
||||
hostVerify(size, A, coef_, &sec_hv);
|
||||
kernelVerify(size, A, coef_, &sec_kv);
|
||||
HIP_CHECK(hipFree(A));
|
||||
|
||||
log_kernel("Managed: kernel fill", GBytes, sec, sec_hv, sec_kv);
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool testManagedMemory() {
|
||||
if (!supportManagedMemory()) {
|
||||
return false;
|
||||
}
|
||||
|
||||
std::cout << "Test managed memory host filling" << std::endl;
|
||||
for (int i = 0; i < NUM_SIZE; i++) {
|
||||
if (!testManagedMemoryHostFill(totalSizes_[i])) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
std::cout << "Test managed memory kernel filling" << std::endl;
|
||||
for (int i = 0; i < NUM_SIZE; i++) {
|
||||
if (!testManagedMemoryKernelFill(totalSizes_[i])) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool testHostMemoryHostFill(size_t size, unsigned int flags) {
|
||||
double GBytes = static_cast<double>(size) / NUM_1GB;
|
||||
T *A;
|
||||
HIP_CHECK(hipHostMalloc(&A, size, flags));
|
||||
double sec = 0;
|
||||
hostFill(size, A, coef_, &sec);
|
||||
HIP_CHECK(hipHostFree(A));
|
||||
|
||||
log_host("Host: host fill", GBytes, sec);
|
||||
return true;
|
||||
}
|
||||
|
||||
bool testHostMemoryKernelFill(size_t size, unsigned int flags) {
|
||||
double GBytes = static_cast<double>(size) / NUM_1GB;
|
||||
|
||||
T *A;
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void **>(&A), size, flags));
|
||||
double sec = 0, sec_hv = 0, sec_kv = 0;
|
||||
kernelFill(size, A, coef_, &sec);
|
||||
hostVerify(size, A, coef_, &sec_hv);
|
||||
kernelVerify(size, A, coef_, &sec_kv);
|
||||
HIP_CHECK(hipHostFree(A));
|
||||
|
||||
log_kernel("Host: kernel fill", GBytes, sec, sec_hv, sec_kv);
|
||||
return true;
|
||||
}
|
||||
|
||||
bool testHostMemory() {
|
||||
std::cout << "Test coherent host memory host filling" << std::endl;
|
||||
for (int i = 0; i < NUM_SIZE; i++) {
|
||||
if (!testHostMemoryHostFill(totalSizes_[i], hipHostMallocCoherent)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
std::cout << "Test non-coherent host memory host filling" << std::endl;
|
||||
for (int i = 0; i < NUM_SIZE; i++) {
|
||||
if (!testHostMemoryHostFill(totalSizes_[i], hipHostMallocNonCoherent)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
std::cout << "Test coherent host memory kernel filling" << std::endl;
|
||||
for (int i = 0; i < NUM_SIZE; i++) {
|
||||
if (!testHostMemoryKernelFill(totalSizes_[i], hipHostMallocCoherent)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
std::cout << "Test non-coherent host memory kernel filling" << std::endl;
|
||||
for (int i = 0; i < NUM_SIZE; i++) {
|
||||
if (!testHostMemoryKernelFill(totalSizes_[i], hipHostMallocNonCoherent)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
/* This function should be via device attribute query*/
|
||||
bool supportDeviceMallocFinegrained() {
|
||||
#ifdef __HIP_PLATFORM_AMD__
|
||||
T *A = nullptr;
|
||||
hipError_t err;
|
||||
err = hipExtMallocWithFlags(reinterpret_cast<void**>(&A), sizeof(T),
|
||||
hipDeviceMallocFinegrained);
|
||||
if (err || !A) {
|
||||
return false;
|
||||
}
|
||||
HIP_CHECK(hipFree(A));
|
||||
return true;
|
||||
#else
|
||||
return false;
|
||||
#endif
|
||||
}
|
||||
|
||||
unsigned int setNumBlocks(size_t size) {
|
||||
size_t num = size/sizeof(T);
|
||||
|
||||
#if USE_HIPTEST_SETNUMBLOCKS
|
||||
return HipTest::setNumBlocks(blocksPerCU_, threadsPerBlock_, num);
|
||||
#else
|
||||
return (num + threadsPerBlock_ - 1) / threadsPerBlock_;
|
||||
#endif
|
||||
}
|
||||
|
||||
#ifdef __HIP_PLATFORM_AMD__
|
||||
bool testExtDeviceMemoryHostFill(size_t size, unsigned int flags) {
|
||||
double GBytes = static_cast<double>(size) / NUM_1GB;
|
||||
|
||||
T *A = nullptr;
|
||||
HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast<void **>(&A),
|
||||
size, flags));
|
||||
if (!A) {
|
||||
std::cout << "failed hipExtMallocWithFlags() with size =" <<
|
||||
size << " flags="<< std::hex << flags << std::endl;
|
||||
return false;
|
||||
}
|
||||
|
||||
double sec = 0;
|
||||
hostFill(size, A, coef_, &sec); // Cpu can access this mem
|
||||
HIP_CHECK(hipFree(A));
|
||||
|
||||
log_host("ExtDevice: host fill", GBytes, sec);
|
||||
return true;
|
||||
}
|
||||
|
||||
bool testExtDeviceMemoryKernelFill(size_t size, unsigned int flags) {
|
||||
double GBytes = static_cast<double>(size) / NUM_1GB;
|
||||
|
||||
T *A = nullptr;
|
||||
HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast<void **>(&A),
|
||||
size, flags));
|
||||
if (!A) {
|
||||
std::cout << "failed hipExtMallocWithFlags() with size =" <<
|
||||
size << " flags=" << std::hex << flags << std::endl;
|
||||
return false;
|
||||
}
|
||||
|
||||
double sec = 0, sec_hv = 0, sec_kv = 0;
|
||||
kernelFill(size, A, coef_, &sec);
|
||||
// Fine grained device memory can be verified by host
|
||||
hostVerify(size, A, coef_, &sec_hv);
|
||||
kernelVerify(size, A, coef_, &sec_kv);
|
||||
HIP_CHECK(hipFree(A));
|
||||
|
||||
log_kernel("ExtDevice: kernel fill", GBytes, sec, sec_hv, sec_kv);
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool testExtDeviceMemory() {
|
||||
std::cout << "Test fine grained device memory host filling"
|
||||
<< std::endl;
|
||||
for (int i = 0; i < NUM_SIZE; i++) {
|
||||
if (!testExtDeviceMemoryHostFill(totalSizes_[i],
|
||||
hipDeviceMallocFinegrained)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
std::cout << "Test fine grained device memory kernel filling"
|
||||
<< std::endl;
|
||||
for (int i = 0; i < NUM_SIZE; i++) {
|
||||
if (!testExtDeviceMemoryKernelFill(totalSizes_[i],
|
||||
hipDeviceMallocFinegrained)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
#endif
|
||||
|
||||
bool run() {
|
||||
if (supportLargeBar()) {
|
||||
if (!testLargeBar()) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
if (supportManagedMemory()) {
|
||||
if (!testManagedMemory()) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
if (!testHostMemory()) {
|
||||
return false;
|
||||
}
|
||||
|
||||
#ifdef __HIP_PLATFORM_AMD__
|
||||
if (supportDeviceMallocFinegrained()) {
|
||||
if (!testExtDeviceMemory()) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
#endif
|
||||
return true;
|
||||
}
|
||||
};
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipPerfMemFill status.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - perftests/memory/hipPerfMemFill.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.6
|
||||
*/
|
||||
|
||||
TEST_CASE("Perf_hipPerfMemFill_test") {
|
||||
std::cout << "Test int" << std::endl;
|
||||
hipPerfMemFill<int> hipPerfMemFillInt;
|
||||
REQUIRE(true == hipPerfMemFillInt.open(0));
|
||||
REQUIRE(true == hipPerfMemFillInt.run());
|
||||
|
||||
std::cout << "Test double" << std::endl;
|
||||
hipPerfMemFill<double> hipPerfMemFillDouble;
|
||||
REQUIRE(true == hipPerfMemFillDouble.open(0));
|
||||
REQUIRE(true == hipPerfMemFillDouble.run());
|
||||
}
|
||||
@@ -0,0 +1,144 @@
|
||||
/*
|
||||
Copyright (c) 2023 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.
|
||||
*/
|
||||
|
||||
/**
|
||||
* @addtogroup hipMemcpy hipMemcpy
|
||||
* @{
|
||||
* @ingroup perfMemoryTest
|
||||
* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` -
|
||||
* Copies data between host and device.
|
||||
*/
|
||||
|
||||
#include <time.h>
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
#define NUM_SIZE 19 // size up to 16M
|
||||
#define NUM_ITER 500 // Total GPU memory up to 16M*500=8G
|
||||
|
||||
void valSet(int* A, int val, size_t size) {
|
||||
size_t len = size / sizeof(int);
|
||||
for (int i = 0; i < len; i++) {
|
||||
A[i] = val;
|
||||
}
|
||||
}
|
||||
|
||||
void setup(size_t *size, int *num, int **pA, const size_t totalGlobalMem) {
|
||||
for (int i = 0; i < *num; i++) {
|
||||
size[i] = 1 << (i + 6);
|
||||
if ((NUM_ITER + 1) * size[i] > totalGlobalMem) {
|
||||
*num = i;
|
||||
break;
|
||||
}
|
||||
}
|
||||
*pA = reinterpret_cast<int*>(malloc(size[*num - 1]));
|
||||
valSet(*pA, 1, size[*num - 1]);
|
||||
}
|
||||
|
||||
void testInit(size_t size, int *A) {
|
||||
int *Ad;
|
||||
|
||||
clock_t start = clock();
|
||||
HIP_CHECK(hipMalloc(&Ad, size)); // hip::init() will be called
|
||||
clock_t end = clock();
|
||||
double uS = (end - start) * 1000000. / CLOCKS_PER_SEC;
|
||||
INFO("Initial: hipMalloc(" << size << ") cost " << uS << "us" << "\n");
|
||||
|
||||
start = clock();
|
||||
HIP_CHECK(hipMemcpy(Ad, A, size, hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
end = clock();
|
||||
uS = (end - start) * 1000000. / CLOCKS_PER_SEC;
|
||||
INFO("hipMemcpy(" << size << ") cost " << uS << "us" << "\n");
|
||||
|
||||
start = clock();
|
||||
HIP_CHECK(hipFree(Ad));
|
||||
end = clock();
|
||||
uS = (end - start) * 1000000. / CLOCKS_PER_SEC;
|
||||
INFO("hipFree(" << size << ") cost " << uS << "us" << "\n");
|
||||
}
|
||||
|
||||
static bool hipPerfMemMallocCpyFree_test() {
|
||||
double uS;
|
||||
clock_t start, end;
|
||||
size_t size[NUM_SIZE] = { 0 };
|
||||
int *Ad[NUM_ITER] = { nullptr };
|
||||
int *A;
|
||||
hipDeviceProp_t props;
|
||||
memset(&props, 0, sizeof(props));
|
||||
HIP_CHECK(hipGetDeviceProperties(&props, 0));
|
||||
INFO("totalGlobalMem: " << props.totalGlobalMem << "\n");
|
||||
|
||||
int num = NUM_SIZE;
|
||||
setup(size, &num, &A, props.totalGlobalMem);
|
||||
testInit(size[0], A);
|
||||
|
||||
for (int i = 0; i < num; i++) {
|
||||
start = clock();
|
||||
for (int j = 0; j < NUM_ITER; j++) {
|
||||
HIP_CHECK(hipMalloc(&Ad[j], size[i]));
|
||||
}
|
||||
end = clock();
|
||||
uS = (end - start) * 1000000. / (NUM_ITER * CLOCKS_PER_SEC);
|
||||
INFO("hipMalloc(" << size[i] << ") cost " << uS << "us" << "\n");
|
||||
|
||||
start = clock();
|
||||
for (int j = 0; j < NUM_ITER; j++) {
|
||||
HIP_CHECK(hipMemcpy(Ad[j], A, size[i], hipMemcpyHostToDevice));
|
||||
}
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
end = clock();
|
||||
uS = (end - start) * 1000000. / (NUM_ITER * CLOCKS_PER_SEC);
|
||||
INFO("hipMemcpy(" << size[i] << ") cost " << uS << "us" << "\n");
|
||||
|
||||
start = clock();
|
||||
for (int j = 0; j < NUM_ITER; j++) {
|
||||
HIP_CHECK(hipFree(Ad[j]));
|
||||
Ad[j] = nullptr;
|
||||
}
|
||||
end = clock();
|
||||
double uS = (end - start) * 1000000. / (NUM_ITER * CLOCKS_PER_SEC);
|
||||
INFO("hipFree(" << size[i] << ") cost " << uS << "us" << "\n");
|
||||
}
|
||||
free(A);
|
||||
return true;
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipPerfMemMallocCpyFree status.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - perftests/memory/hipPerfMemMallocCpyFree.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.6
|
||||
*/
|
||||
|
||||
TEST_CASE("Perf_hipPerfMemMallocCpyFree_test") {
|
||||
int numDevices = 0;
|
||||
HIP_CHECK(hipGetDeviceCount(&numDevices));
|
||||
|
||||
if (numDevices <= 0) {
|
||||
SUCCEED("Skipped testcase hipPerfDevMemReadSpeed as"
|
||||
"there is no device to test.");
|
||||
} else {
|
||||
REQUIRE(true == hipPerfMemMallocCpyFree_test());
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,118 @@
|
||||
/*
|
||||
Copyright (c) 2023 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.
|
||||
*/
|
||||
|
||||
/**
|
||||
* @addtogroup hipMemcpy hipMemcpy
|
||||
* @{
|
||||
* @ingroup perfMemoryTest
|
||||
* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` -
|
||||
* Copies data between host and device.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
#define NUM_SIZE 8
|
||||
#define NUM_ITER 0x40000
|
||||
|
||||
class hipPerfMemcpy {
|
||||
private:
|
||||
unsigned int numBuffers_;
|
||||
size_t totalSizes_[NUM_SIZE];
|
||||
void setHostBuffer(int *A, int val, size_t size);
|
||||
public:
|
||||
hipPerfMemcpy();
|
||||
~hipPerfMemcpy() {}
|
||||
bool run(unsigned int numTests);
|
||||
};
|
||||
|
||||
hipPerfMemcpy::hipPerfMemcpy() : numBuffers_(0) {
|
||||
for (int i = 0; i < NUM_SIZE; i++) {
|
||||
totalSizes_[i] = 1 << (i + 6);
|
||||
}
|
||||
}
|
||||
|
||||
void hipPerfMemcpy::setHostBuffer(int *A, int val, size_t size) {
|
||||
size_t len = size / sizeof(int);
|
||||
for (int i = 0; i < len; i++) {
|
||||
A[i] = val;
|
||||
}
|
||||
}
|
||||
|
||||
bool hipPerfMemcpy::run(unsigned int numTests) {
|
||||
int *A, *Ad;
|
||||
A = new int[totalSizes_[numTests]];
|
||||
setHostBuffer(A, 1, totalSizes_[numTests]);
|
||||
HIP_CHECK(hipMalloc(&Ad, totalSizes_[numTests]));
|
||||
|
||||
// measure performance based on host time
|
||||
auto all_start = std::chrono::steady_clock::now();
|
||||
|
||||
for (int j = 0; j < NUM_ITER; j++) {
|
||||
HIP_CHECK(hipMemcpy(Ad, A, totalSizes_[numTests], hipMemcpyHostToDevice));
|
||||
}
|
||||
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
auto all_end = std::chrono::steady_clock::now();
|
||||
std::chrono::duration<double, std::micro> diff = all_end - all_start;
|
||||
|
||||
INFO("hipPerfMemcpy[" << numTests << "] " << "Host to Device copy took "
|
||||
<< diff.count() / NUM_ITER << " sec for memory size of " <<
|
||||
totalSizes_[numTests] << " Bytes.");
|
||||
|
||||
delete [] A;
|
||||
HIP_CHECK(hipFree(Ad));
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipPerfMemcpy status.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - perftests/memory/hipPerfMemcpy.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.6
|
||||
*/
|
||||
|
||||
TEST_CASE("Perf_hipPerfMemcpy_test") {
|
||||
int numDevices = 0;
|
||||
HIP_CHECK(hipGetDeviceCount(&numDevices));
|
||||
|
||||
if (numDevices <= 0) {
|
||||
SUCCEED("Skipped testcase hipPerfMemcpy as there is no device to test.");
|
||||
} else {
|
||||
int deviceId = 0;
|
||||
HIP_CHECK(hipSetDevice(deviceId));
|
||||
hipDeviceProp_t props = {0};
|
||||
HIP_CHECK(hipGetDeviceProperties(&props, deviceId));
|
||||
|
||||
INFO("info: running on bus " << "0x" << props.pciBusID << " " <<
|
||||
props.name << " with " << props.multiProcessorCount << " CUs "
|
||||
<< " and device id: " << deviceId);
|
||||
|
||||
hipPerfMemcpy hipPerfMemcpy;
|
||||
for (auto testCase = 0; testCase < NUM_SIZE; testCase++) {
|
||||
REQUIRE(true == hipPerfMemcpy.run(testCase));
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,419 @@
|
||||
/*
|
||||
Copyright (c) 2023 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.
|
||||
*/
|
||||
|
||||
/**
|
||||
* @addtogroup hipMemsetKernel hipMemsetKernel
|
||||
* @{
|
||||
* @ingroup perfMemoryTest
|
||||
* `hipMemset(void* devPtr, int value, size_t count)` -
|
||||
* Initializes or sets device memory to a value.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
static size_t typeSizeList[] = {
|
||||
1, 2, 4, 8, 16, 32, 64, 128,
|
||||
};
|
||||
|
||||
static unsigned int sizeList[] = {
|
||||
256, 512, 1024, 2048, 4096, 8192,
|
||||
};
|
||||
|
||||
static unsigned int eleNumList[] = {
|
||||
0x0020000, 0x0080000, 0x0200000, 0x0800000, 0x2000000,
|
||||
};
|
||||
|
||||
typedef struct _dataType {
|
||||
char memsetval = 0x42;
|
||||
char memsetD8val = 0xDE;
|
||||
int16_t memsetD16val = 0xDEAD;
|
||||
int memsetD32val = 0xDEADBEEF;
|
||||
}dataType;
|
||||
|
||||
#define NUM_ITER 100
|
||||
|
||||
enum MemsetType {
|
||||
hipMemsetTypeDefault,
|
||||
hipMemsetTypeD8,
|
||||
hipMemsetTypeD16,
|
||||
hipMemsetTypeD32
|
||||
};
|
||||
|
||||
class hipPerfMemset {
|
||||
private:
|
||||
unsigned int bufSize_;
|
||||
unsigned int num_typeSize_;
|
||||
unsigned int num_elements_;
|
||||
size_t testTypeSize_;
|
||||
unsigned int testNumEle_;
|
||||
unsigned int _numSubTests = 0;
|
||||
unsigned int _numSubTests2D = 0;
|
||||
unsigned int _numSubTests3D = 0;
|
||||
unsigned int num_sizes_ = 0;
|
||||
|
||||
public:
|
||||
hipPerfMemset() {
|
||||
num_typeSize_ = sizeof(typeSizeList) / sizeof(size_t);
|
||||
num_elements_ = sizeof(eleNumList) / sizeof(unsigned int);
|
||||
_numSubTests = num_elements_ * num_typeSize_;
|
||||
|
||||
num_sizes_ = sizeof(sizeList) / sizeof(unsigned int);
|
||||
_numSubTests2D = num_sizes_;
|
||||
_numSubTests3D = _numSubTests2D;
|
||||
}
|
||||
|
||||
~hipPerfMemset() {}
|
||||
|
||||
bool open(int deviceID);
|
||||
|
||||
template<typename T>
|
||||
void run1D(unsigned int test, T memsetval, enum MemsetType type, bool async);
|
||||
|
||||
template<typename T>
|
||||
void run2D(unsigned int test, T memsetval, enum MemsetType type, bool async);
|
||||
|
||||
template<typename T>
|
||||
void run3D(unsigned int test, T memsetval, enum MemsetType type, bool async);
|
||||
|
||||
uint getNumTests() {
|
||||
return _numSubTests;
|
||||
}
|
||||
|
||||
uint getNumTests2D() {
|
||||
return _numSubTests2D;
|
||||
}
|
||||
uint getNumTests3D() {
|
||||
return _numSubTests3D;
|
||||
}
|
||||
};
|
||||
|
||||
bool hipPerfMemset::open(int deviceId) {
|
||||
int nGpu = 0;
|
||||
HIP_CHECK(hipGetDeviceCount(&nGpu));
|
||||
if (nGpu < 1) {
|
||||
return false;
|
||||
}
|
||||
|
||||
HIP_CHECK(hipSetDevice(deviceId));
|
||||
hipDeviceProp_t props = {0};
|
||||
HIP_CHECK(hipGetDeviceProperties(&props, deviceId));
|
||||
INFO("info: running on bus " << "0x" << props.pciBusID << " " << props.name
|
||||
<< " with " << props.multiProcessorCount << " CUs and device id: "
|
||||
<< deviceId << "\n");
|
||||
return true;
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void hipPerfMemset::run1D(unsigned int test, T memsetval,
|
||||
enum MemsetType type, bool async) {
|
||||
T *A_h, *A_d;
|
||||
|
||||
testTypeSize_ = typeSizeList[(test / num_elements_) % num_typeSize_];
|
||||
testNumEle_ = eleNumList[test % num_elements_];
|
||||
|
||||
bufSize_ = testNumEle_ * 4;
|
||||
|
||||
HIP_CHECK(hipMalloc(&A_d, bufSize_));
|
||||
|
||||
A_h = reinterpret_cast<T*> (malloc(bufSize_));
|
||||
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
// Warm-up
|
||||
HIP_CHECK(hipMemset(reinterpret_cast<void *>(A_d), memsetval, bufSize_));
|
||||
|
||||
auto start = std::chrono::steady_clock::now();
|
||||
|
||||
for (uint i = 0; i < NUM_ITER; i++) {
|
||||
if (type == hipMemsetTypeDefault && !async) {
|
||||
HIP_CHECK(hipMemset(reinterpret_cast<void *>(A_d), memsetval, bufSize_));
|
||||
} else if (type == hipMemsetTypeDefault && async) {
|
||||
HIP_CHECK(hipMemsetAsync(A_d, memsetval, bufSize_, stream));
|
||||
} else if (type == hipMemsetTypeD8 && !async) {
|
||||
HIP_CHECK(hipMemsetD8((hipDeviceptr_t)A_d, memsetval, bufSize_));
|
||||
} else if (type == hipMemsetTypeD8 && async) {
|
||||
HIP_CHECK(hipMemsetD8Async((hipDeviceptr_t)A_d, memsetval, bufSize_));
|
||||
} else if (type == hipMemsetTypeD16 && !async) {
|
||||
HIP_CHECK(hipMemsetD16((hipDeviceptr_t)A_d, memsetval,
|
||||
bufSize_/sizeof(T)));
|
||||
} else if (type == hipMemsetTypeD16 && async) {
|
||||
HIP_CHECK(hipMemsetD16Async((hipDeviceptr_t)A_d, memsetval,
|
||||
bufSize_/sizeof(T)));
|
||||
} else if (type == hipMemsetTypeD32 && !async) {
|
||||
HIP_CHECK(hipMemsetD32((hipDeviceptr_t)A_d, memsetval,
|
||||
bufSize_/sizeof(T)));
|
||||
} else if (type == hipMemsetTypeD32 && async) {
|
||||
HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)A_d, memsetval,
|
||||
bufSize_/sizeof(T)));
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
auto end = std::chrono::steady_clock::now();
|
||||
|
||||
HIP_CHECK(hipMemcpy(A_h, A_d, bufSize_, hipMemcpyDeviceToHost) );
|
||||
|
||||
for (int i = 0; i < bufSize_/testTypeSize_; i++) {
|
||||
if (A_h[i] != memsetval) {
|
||||
INFO("mismatch at index " << i << " computed: " <<
|
||||
static_cast<int> (A_h[i]) << ", memsetval: " <<
|
||||
static_cast<int> (memsetval) << "\n");
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(A_d));
|
||||
free(A_h);
|
||||
|
||||
std::chrono::duration<double> diff = end - start;
|
||||
|
||||
auto sec = diff.count();
|
||||
auto perf = static_cast<double>((bufSize_ * NUM_ITER * (1e-09)) / sec);
|
||||
|
||||
INFO("hipPerf1DMemset[" << test << "] " << (int)bufSize_/1024 << " Kb "
|
||||
<< std::setw(4) << " typeSize " << (int) testTypeSize_ << ":"
|
||||
<< std::setw(5) << perf << " GB/s \n");
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void hipPerfMemset::run2D(unsigned int test, T memsetval,
|
||||
enum MemsetType type, bool async) {
|
||||
bufSize_ = sizeList[test % num_sizes_];
|
||||
size_t numH = bufSize_;
|
||||
size_t numW = bufSize_;
|
||||
size_t pitch_A;
|
||||
size_t width = numW * sizeof(char);
|
||||
size_t sizeElements = width * numH;
|
||||
size_t elements = numW* numH;
|
||||
|
||||
T * A_h, * A_d;
|
||||
|
||||
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d),
|
||||
&pitch_A, width, numH));
|
||||
A_h = reinterpret_cast<char*>(malloc(sizeElements));
|
||||
|
||||
for (size_t i=0; i < elements; i++) {
|
||||
A_h[i] = 1;
|
||||
}
|
||||
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
// Warm-up
|
||||
HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, numW, numH));
|
||||
|
||||
auto start = std::chrono::steady_clock::now();
|
||||
|
||||
for (uint i = 0; i < NUM_ITER; i++) {
|
||||
if (type == hipMemsetTypeDefault && !async) {
|
||||
HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, numW, numH));
|
||||
} else if (type == hipMemsetTypeDefault && async) {
|
||||
HIP_CHECK(hipMemset2DAsync(A_d, pitch_A, memsetval, numW, numH, stream));
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
|
||||
auto end = std::chrono::steady_clock::now();
|
||||
|
||||
HIP_CHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH,
|
||||
hipMemcpyDeviceToHost));
|
||||
|
||||
for (int i=0; i < elements; i++) {
|
||||
if (A_h[i] != memsetval) {
|
||||
INFO("mismatch at index " << i << " computed: " <<
|
||||
static_cast<int> (A_h[i]) << ", memsetval: " <<
|
||||
static_cast<int> (memsetval) << "\n");
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
|
||||
std::chrono::duration<double> diff = end - start;
|
||||
|
||||
auto sec = diff.count();
|
||||
auto perf = static_cast<double>((sizeElements* NUM_ITER * (1e-09)) / sec);
|
||||
|
||||
INFO("hipPerf2DMemset[" << test << "] " <<" " << "(GB/s) for " <<
|
||||
(int)bufSize_ << " x " << bufSize_ << " bytes : " << std::setw(5) <<
|
||||
perf << "\n");
|
||||
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
HIP_CHECK(hipFree(A_d));
|
||||
free(A_h);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void hipPerfMemset::run3D(unsigned int test, T memsetval,
|
||||
enum MemsetType type, bool async) {
|
||||
bufSize_ = sizeList[test % num_sizes_];
|
||||
|
||||
size_t numH = bufSize_;
|
||||
size_t numW = bufSize_;
|
||||
size_t depth = 10;
|
||||
size_t width = numW * sizeof(char);
|
||||
size_t sizeElements = width * numH * depth;
|
||||
size_t elements = numW* numH* depth;
|
||||
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
T *A_h;
|
||||
|
||||
hipExtent extent = make_hipExtent(width, numH, depth);
|
||||
hipPitchedPtr devPitchedPtr;
|
||||
|
||||
HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
|
||||
A_h = reinterpret_cast<char*>(malloc(sizeElements));
|
||||
HIPASSERT(A_h != NULL);
|
||||
|
||||
for (size_t i = 0; i < elements; i++) {
|
||||
A_h[i] = 1;
|
||||
}
|
||||
|
||||
// Warm up
|
||||
HIP_CHECK(hipMemset3D(devPitchedPtr, memsetval, extent));
|
||||
|
||||
auto start = std::chrono::steady_clock::now();
|
||||
|
||||
for (uint i = 0; i < NUM_ITER; i++) {
|
||||
if (type == hipMemsetTypeDefault && !async) {
|
||||
HIP_CHECK(hipMemset3D(devPitchedPtr, memsetval, extent));
|
||||
} else if (type == hipMemsetTypeDefault && async) {
|
||||
HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream));
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamSynchronize(stream));
|
||||
|
||||
auto end = std::chrono::steady_clock::now();
|
||||
|
||||
hipMemcpy3DParms myparms = {0};
|
||||
myparms.srcPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPos = make_hipPos(0, 0, 0);
|
||||
myparms.dstPtr = make_hipPitchedPtr(A_h, width , numW, numH);
|
||||
myparms.srcPtr = devPitchedPtr;
|
||||
myparms.extent = extent;
|
||||
|
||||
myparms.kind = hipMemcpyDeviceToHost;
|
||||
|
||||
HIP_CHECK(hipMemcpy3D(&myparms));
|
||||
|
||||
for (int i=0; i < elements; i++) {
|
||||
if (A_h[i] != memsetval) {
|
||||
INFO("mismatch at index " << i << " computed: " <<
|
||||
static_cast<int> (A_h[i]) << ", memsetval: " <<
|
||||
static_cast<int> (memsetval) << "\n");
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
|
||||
std::chrono::duration<double> diff = end - start;
|
||||
|
||||
auto sec = diff.count();
|
||||
auto perf = static_cast<double>((sizeElements * NUM_ITER * (1e-09)) / sec);
|
||||
|
||||
INFO("hipPerf3DMemset[" << test << "] " <<" " << "(GB/s) for " <<
|
||||
(int)bufSize_ << " x " << bufSize_ << " x " <<depth << " bytes : " <<
|
||||
std::setw(5) << perf << "\n");
|
||||
HIP_CHECK(hipFree(devPitchedPtr.ptr));
|
||||
free(A_h);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipPerfMemset status.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - perftests/memory/hipPerfMemset.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.6
|
||||
*/
|
||||
|
||||
TEST_CASE("Perf_hipPerfMemset_test") {
|
||||
hipPerfMemset hipPerfMemset;
|
||||
|
||||
int deviceId = 0;
|
||||
REQUIRE(hipPerfMemset.open(deviceId));
|
||||
|
||||
dataType pattern;
|
||||
|
||||
int numTests = hipPerfMemset.getNumTests();
|
||||
int numTests2D = hipPerfMemset.getNumTests2D();
|
||||
int numTests3D = hipPerfMemset.getNumTests3D();
|
||||
|
||||
bool async = false;
|
||||
|
||||
for (uint i = 0; i < 2 ; i++) {
|
||||
if (async) {
|
||||
INFO("Perf of hipMemsetAsync for 1D arrays \n");
|
||||
} else {
|
||||
INFO("Perf of hipMemset for 1D arrays \n");
|
||||
}
|
||||
|
||||
for (auto testCase = 0; testCase < numTests; testCase++) {
|
||||
if (testCase < 5) {
|
||||
INFO("API: hipMemset \n");
|
||||
hipPerfMemset.run1D(testCase, pattern.memsetval,
|
||||
hipMemsetTypeDefault, async);
|
||||
} else if (testCase < 10) {
|
||||
INFO("API: hipMemsetD16 \n");
|
||||
hipPerfMemset.run1D(testCase, pattern.memsetD16val,
|
||||
hipMemsetTypeD16, async);
|
||||
} else if (testCase < 15) {
|
||||
INFO("API: hipMemsetD32 \n");
|
||||
hipPerfMemset.run1D(testCase, pattern.memsetD32val,
|
||||
hipMemsetTypeD32, async);
|
||||
} else {
|
||||
INFO("API: hipMemset \n");
|
||||
hipPerfMemset.run1D(testCase, pattern.memsetval,
|
||||
hipMemsetTypeDefault, async);
|
||||
}
|
||||
}
|
||||
async = true;
|
||||
}
|
||||
|
||||
for (uint i = 0; i < 2; i++) {
|
||||
if (async) {
|
||||
INFO("Perf of hipMemset2DAsync for 2D arrays \n");
|
||||
} else {
|
||||
INFO("Perf of hipMemset2D for 2D arrays \n");
|
||||
}
|
||||
|
||||
for (uint test = 0; test < numTests2D; test++) {
|
||||
hipPerfMemset.run2D(test, pattern.memsetval, hipMemsetTypeDefault, async);
|
||||
}
|
||||
async = false;
|
||||
}
|
||||
|
||||
for (uint i = 0; i < 2; i++) {
|
||||
if (async) {
|
||||
INFO("Perf of hipMemset3DAsync for 3D arrays \n");
|
||||
} else {
|
||||
INFO("Perf of hipMemset3D for 3D arrays \n");
|
||||
}
|
||||
|
||||
for (uint test = 0; test < numTests3D; test++) {
|
||||
hipPerfMemset.run3D(test, pattern.memsetval, hipMemsetTypeDefault, async);
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,311 @@
|
||||
/*
|
||||
Copyright (c) 2023 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.
|
||||
*/
|
||||
|
||||
|
||||
/**
|
||||
* @addtogroup hipMemcpyKernel hipMemcpyKernel
|
||||
* @{
|
||||
* @ingroup perfMemoryTest
|
||||
* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` -
|
||||
* Copies data between host and device.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
#define NUM_TYPES 3
|
||||
std::vector<std::string> types = {"float", "float2", "float4"};
|
||||
std::vector<unsigned int> typeSizes = {4, 8, 16};
|
||||
|
||||
#define NUM_SIZES 12
|
||||
std::vector<unsigned int> sizes = {1, 2, 4, 8, 16, 32,
|
||||
64, 128, 256, 512, 1024, 2048};
|
||||
|
||||
#define NUM_BUFS 6
|
||||
#define MAX_BUFS (1 << (NUM_BUFS - 1))
|
||||
|
||||
#ifdef __HIP_PLATFORM_NVIDIA__
|
||||
__host__ __device__ void operator+=(float2 &a, float2 b) { //NOLINT
|
||||
a.x += b.x; a.y += b.y;
|
||||
}
|
||||
|
||||
__host__ __device__ void operator+=(float4 &a, float4 b) { //NOLINT
|
||||
a.x += b.x; a.y += b.y; a.z += b.z; a.w += b.w;
|
||||
}
|
||||
#endif
|
||||
|
||||
template <typename T>
|
||||
__global__ void sampleRate(T * outBuffer, unsigned int inBufSize,
|
||||
unsigned int writeIt, T **inBuffer, int numBufs) {
|
||||
uint gid = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
uint inputIdx = gid % inBufSize;
|
||||
|
||||
T tmp;
|
||||
memset(&tmp, 0, sizeof(T));
|
||||
for (int i = 0; i < numBufs; i++) {
|
||||
tmp += *(*(inBuffer+i)+inputIdx);
|
||||
}
|
||||
|
||||
if (writeIt*(unsigned int)tmp.x) {
|
||||
outBuffer[gid] = tmp;
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
__global__ void sampleRateFloat(T * outBuffer, unsigned int inBufSize,
|
||||
unsigned int writeIt, T ** inBuffer, int numBufs) {
|
||||
uint gid = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
uint inputIdx = gid % inBufSize;
|
||||
|
||||
T tmp = (T)0.0f;
|
||||
|
||||
for (int i = 0; i < numBufs; i++) {
|
||||
tmp += *((*inBuffer+i)+inputIdx);
|
||||
}
|
||||
|
||||
if (writeIt*(unsigned int)tmp) {
|
||||
outBuffer[gid] = tmp;
|
||||
}
|
||||
}
|
||||
|
||||
class hipPerfSampleRate {
|
||||
public:
|
||||
hipPerfSampleRate();
|
||||
~hipPerfSampleRate();
|
||||
|
||||
bool open(void);
|
||||
void run(unsigned int testCase);
|
||||
void close(void);
|
||||
|
||||
// array of funtion pointers
|
||||
typedef void (hipPerfSampleRate::*funPtr)(void * outBuffer, unsigned int
|
||||
inBufSize, unsigned int writeIt, void **inBuffer, int numBufs,
|
||||
int grids, int blocks, int threads_per_block);
|
||||
|
||||
// Wrappers
|
||||
void float_kernel(void * outBuffer, unsigned int inBufSize,
|
||||
unsigned int writeIt, void **inBuffer, int numBufs,
|
||||
int grids, int blocks, int threads_per_block);
|
||||
|
||||
void float2_kernel(void * outBuffer, unsigned int inBufSize,
|
||||
unsigned int writeIt, void **inBuffer, int numBufs,
|
||||
int grids, int blocks, int threads_per_block);
|
||||
|
||||
void float4_kernel(void * outBuffer, unsigned int inBufSize,
|
||||
unsigned int writeIt, void **inBuffer, int numBufs,
|
||||
int grids, int blocks, int threads_per_block);
|
||||
|
||||
private:
|
||||
void setData(void *ptr, unsigned int value);
|
||||
void checkData(uint *ptr);
|
||||
|
||||
unsigned int width_;
|
||||
unsigned int bufSize_;
|
||||
int numCUs;
|
||||
|
||||
unsigned int outBufSize_;
|
||||
static const unsigned int MAX_ITERATIONS = 25;
|
||||
unsigned int numBufs_;
|
||||
unsigned int typeIdx_;
|
||||
};
|
||||
|
||||
hipPerfSampleRate::hipPerfSampleRate() {}
|
||||
hipPerfSampleRate::~hipPerfSampleRate() {}
|
||||
void hipPerfSampleRate::close() {}
|
||||
|
||||
bool hipPerfSampleRate::open(void) {
|
||||
int nGpu = 0;
|
||||
HIP_CHECK(hipGetDeviceCount(&nGpu));
|
||||
if (nGpu < 1) {
|
||||
return false;
|
||||
}
|
||||
|
||||
int deviceId = 0;
|
||||
hipDeviceProp_t props = {0};
|
||||
props = {0};
|
||||
HIP_CHECK(hipSetDevice(deviceId));
|
||||
HIP_CHECK(hipGetDeviceProperties(&props, deviceId));
|
||||
INFO("info: running on bus " << "0x" << props.pciBusID << " " <<
|
||||
props.name << " with " << props.multiProcessorCount <<
|
||||
" CUs" << " and device id: " << deviceId << "\n");
|
||||
numCUs = props.multiProcessorCount;
|
||||
return true;
|
||||
}
|
||||
|
||||
// Wrappers for the kernel launches
|
||||
void hipPerfSampleRate::float_kernel(void * outBuffer, unsigned int inBufSize,
|
||||
unsigned int writeIt, void **inBuffer, int numBufs,
|
||||
int grids, int blocks, int threads_per_block) {
|
||||
hipLaunchKernelGGL(sampleRateFloat<float>, dim3(grids, grids, grids),
|
||||
dim3(blocks), 0, 0, reinterpret_cast<float*>(outBuffer),
|
||||
inBufSize, writeIt, reinterpret_cast<float**>(inBuffer), numBufs);
|
||||
}
|
||||
|
||||
void hipPerfSampleRate::float2_kernel(void * outBuffer, unsigned int inBufSize,
|
||||
unsigned int writeIt, void **inBuffer, int grids,
|
||||
int blocks, int threads_per_block, int numBufs) {
|
||||
hipLaunchKernelGGL(sampleRate<float2>, dim3(grids, grids, grids),
|
||||
dim3(blocks), 0, 0, reinterpret_cast<float2 *>(outBuffer),
|
||||
inBufSize, writeIt, reinterpret_cast<float2 **>(inBuffer), numBufs);
|
||||
}
|
||||
|
||||
void hipPerfSampleRate::float4_kernel(void * outBuffer, unsigned int inBufSize,
|
||||
unsigned int writeIt, void **inBuffer, int grids,
|
||||
int blocks, int threads_per_block, int numBufs) {
|
||||
hipLaunchKernelGGL(sampleRate<float4>, dim3(grids, grids, grids),
|
||||
dim3(blocks), 0, 0, reinterpret_cast<float4 *>(outBuffer),
|
||||
inBufSize, writeIt, reinterpret_cast<float4 **>(inBuffer), numBufs);
|
||||
}
|
||||
|
||||
void hipPerfSampleRate::run(unsigned int test) {
|
||||
funPtr p[] = {&hipPerfSampleRate::float_kernel,
|
||||
&hipPerfSampleRate::float2_kernel,
|
||||
&hipPerfSampleRate::float4_kernel};
|
||||
|
||||
// We compute a square domain
|
||||
width_ = sizes[test % NUM_SIZES];
|
||||
typeIdx_ = (test / NUM_SIZES) % NUM_TYPES;
|
||||
bufSize_ = width_ * width_ * typeSizes[typeIdx_];
|
||||
numBufs_ = (1 << (test / (NUM_SIZES * NUM_TYPES)));
|
||||
|
||||
void ** dPtr;
|
||||
void * hOutPtr;
|
||||
void * dOutPtr;
|
||||
void * hInPtr[numBufs_];
|
||||
void * dInPtr[numBufs_];
|
||||
|
||||
outBufSize_ =
|
||||
sizes[NUM_SIZES - 1] * sizes[NUM_SIZES - 1] * typeSizes[NUM_TYPES - 1];
|
||||
|
||||
// Allocate memory on the host and device
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void **>(&hOutPtr), outBufSize_,
|
||||
hipHostMallocDefault));
|
||||
setData(reinterpret_cast<void *>(hOutPtr), 0xdeadbeef);
|
||||
HIP_CHECK(hipMalloc(reinterpret_cast<uint **>(&dOutPtr), outBufSize_));
|
||||
|
||||
// Allocate 2D array in Device
|
||||
HIP_CHECK(hipMalloc(reinterpret_cast<void **>(&dPtr),
|
||||
numBufs_* sizeof(void *)));
|
||||
|
||||
for (uint i = 0; i < numBufs_; i++) {
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void **>(&hInPtr[i]), bufSize_,
|
||||
hipHostMallocDefault));
|
||||
HIP_CHECK(hipMalloc(reinterpret_cast<uint **>(&dInPtr[i]), bufSize_));
|
||||
setData(hInPtr[i], 0x3f800000);
|
||||
}
|
||||
|
||||
// Populate array of pointers with array addresses
|
||||
HIP_CHECK(hipMemcpy(dPtr, dInPtr, numBufs_* sizeof(void *),
|
||||
hipMemcpyHostToDevice));
|
||||
|
||||
// Copy memory from host to device
|
||||
for (uint i = 0; i < numBufs_; i++) {
|
||||
HIP_CHECK(hipMemcpy(dInPtr[i], hInPtr[i], bufSize_, hipMemcpyHostToDevice));
|
||||
}
|
||||
|
||||
HIP_CHECK(hipMemcpy(dOutPtr, hOutPtr, outBufSize_, hipMemcpyHostToDevice));
|
||||
|
||||
// Prepare kernel launch parameters
|
||||
// outBufSize_/sizeof(uint) - Grid size in 3D
|
||||
int grids = 64;
|
||||
int blocks = 64;
|
||||
int threads_per_block = 1;
|
||||
|
||||
unsigned int maxIter = MAX_ITERATIONS * (MAX_BUFS / numBufs_);
|
||||
unsigned int sizeDW = width_ * width_;
|
||||
unsigned int writeIt = 0;
|
||||
|
||||
int idx = 0;
|
||||
|
||||
if (!types[typeIdx_].compare("float")) {
|
||||
idx = 0;
|
||||
} else if (!types[typeIdx_].compare("float2")) {
|
||||
idx = 1;
|
||||
} else if (!types[typeIdx_].compare("float4")) {
|
||||
idx = 2;
|
||||
}
|
||||
|
||||
// Time the kernel execution
|
||||
auto all_start = std::chrono::steady_clock::now();
|
||||
for (uint i = 0; i < maxIter; i++) {
|
||||
(this->*p[idx]) (reinterpret_cast<void *>(dOutPtr), sizeDW, writeIt,
|
||||
dPtr, numBufs_, grids, blocks, threads_per_block);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
auto all_end = std::chrono::steady_clock::now();
|
||||
std::chrono::duration<double> all_kernel_time = all_end - all_start;
|
||||
|
||||
double perf = (static_cast<double>(outBufSize_ * numBufs_ *
|
||||
maxIter * (1e-09))) / all_kernel_time.count();
|
||||
|
||||
INFO("Domain " << sizes[NUM_SIZES - 1] << "x"<< sizes[NUM_SIZES - 1]
|
||||
<< " bufs " << numBufs_ << " " << types[typeIdx_] << " " << width_
|
||||
<< "x" <<width_<< " (GB/s) " << perf << "\n");
|
||||
|
||||
HIP_CHECK(hipFree(dOutPtr));
|
||||
|
||||
// Free host and device memory
|
||||
for (uint i = 0; i < numBufs_; i++) {
|
||||
HIP_CHECK(hipHostFree(hInPtr[i]));
|
||||
HIP_CHECK(hipFree(dInPtr[i]));
|
||||
}
|
||||
HIP_CHECK(hipHostFree(hOutPtr));
|
||||
HIP_CHECK(hipFree(dPtr));
|
||||
}
|
||||
|
||||
|
||||
void hipPerfSampleRate::setData(void *ptr, unsigned int value) {
|
||||
unsigned int *ptr2 = (unsigned int *)ptr;
|
||||
for (unsigned int i = 0; i < bufSize_ / sizeof(unsigned int); i++) {
|
||||
ptr2[i] = value;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
void hipPerfSampleRate::checkData(uint *ptr) {
|
||||
for (unsigned int i = 0; i < outBufSize_ / sizeof(float); i++) {
|
||||
if (ptr[i] != static_cast<float>(numBufs_)) {
|
||||
INFO("Data validation failed at "<< i << " Got "<< ptr[i]
|
||||
<< ", expected " << (float)numBufs_ << "\n");
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipPerfSampleRate status.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - perftests/memory/hipPerfSampleRate.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.6
|
||||
*/
|
||||
|
||||
TEST_CASE("Perf_hipPerfSampleRate_test") {
|
||||
hipPerfSampleRate sampleTypes;
|
||||
|
||||
REQUIRE(true == sampleTypes.open());
|
||||
|
||||
for (unsigned int testCase = 0; testCase < 216 ; testCase+=36) {
|
||||
sampleTypes.run(testCase);
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,262 @@
|
||||
/*
|
||||
Copyright (c) 2023 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.
|
||||
*/
|
||||
|
||||
/**
|
||||
* @addtogroup hipMemcpyKernel hipMemcpyKernel
|
||||
* @{
|
||||
* @ingroup perfMemoryTest
|
||||
* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` -
|
||||
* Copies data between host and device.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
#define sharedMemSize1 2048
|
||||
#define sharedMemSize2 256
|
||||
|
||||
__global__ void sharedMemReadSpeed1(float *outBuf, ulong N) {
|
||||
size_t gid = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
size_t lid = threadIdx.x;
|
||||
__shared__ float local[sharedMemSize1];
|
||||
|
||||
float val1 = 0;
|
||||
float val2 = 0;
|
||||
float val3 = 0;
|
||||
float val4 = 0;
|
||||
|
||||
for (int i = 0; i < (sharedMemSize1 / 64); i++) {
|
||||
local[lid + i * 64] = lid;
|
||||
}
|
||||
|
||||
__syncthreads();
|
||||
|
||||
val1 += local[lid];
|
||||
val2 += local[lid + 64];
|
||||
val3 += local[lid + 128];
|
||||
val4 += local[lid + 192];
|
||||
val1 += local[lid + 256];
|
||||
val2 += local[lid + 320];
|
||||
val3 += local[lid + 384];
|
||||
val4 += local[lid + 448];
|
||||
val1 += local[lid + 512];
|
||||
val2 += local[lid + 576];
|
||||
val3 += local[lid + 640];
|
||||
val4 += local[lid + 704];
|
||||
val1 += local[lid + 768];
|
||||
val2 += local[lid + 832];
|
||||
val3 += local[lid + 896];
|
||||
val4 += local[lid + 960];
|
||||
val1 += local[lid + 1024];
|
||||
val2 += local[lid + 1088];
|
||||
val3 += local[lid + 1152];
|
||||
val4 += local[lid + 1216];
|
||||
val1 += local[lid + 1280];
|
||||
val2 += local[lid + 1344];
|
||||
val3 += local[lid + 1408];
|
||||
val4 += local[lid + 1472];
|
||||
val1 += local[lid + 1536];
|
||||
val2 += local[lid + 1600];
|
||||
val3 += local[lid + 1664];
|
||||
val4 += local[lid + 1728];
|
||||
val1 += local[lid + 1792];
|
||||
val2 += local[lid + 1856];
|
||||
val3 += local[lid + 1920];
|
||||
val4 += local[lid + 1984];
|
||||
|
||||
if (gid < N) {
|
||||
outBuf[gid] = val1 + val2 + val3 + val4;
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void sharedMemReadSpeed2(float *outBuf, ulong N) {
|
||||
size_t gid = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
size_t lid = threadIdx.x;
|
||||
__shared__ float local[sharedMemSize2];
|
||||
|
||||
float val0 = 0.0f;
|
||||
float val1 = 0.0f;
|
||||
|
||||
for (int i = 0; i < (sharedMemSize2 / 64); i++) {
|
||||
local[lid + i * 64] = lid;
|
||||
}
|
||||
|
||||
__syncthreads();
|
||||
|
||||
#pragma nounroll
|
||||
for (uint i = 0; i < 32; i++) {
|
||||
val0 += local[8 * i + 0];
|
||||
val1 += local[8 * i + 1];
|
||||
val0 += local[8 * i + 2];
|
||||
val1 += local[8 * i + 3];
|
||||
val0 += local[8 * i + 4];
|
||||
val1 += local[8 * i + 5];
|
||||
val0 += local[8 * i + 6];
|
||||
val1 += local[8 * i + 7];
|
||||
}
|
||||
|
||||
if (gid < N) {
|
||||
outBuf[gid] = val0 + val1;
|
||||
}
|
||||
}
|
||||
|
||||
static bool hipPerfSharedMemReadSpeed_test() {
|
||||
float *dDst;
|
||||
float *hDst;
|
||||
hipStream_t stream;
|
||||
constexpr uint numSizes = 4;
|
||||
constexpr uint Sizes[numSizes] = {262144, 1048576, 4194304, 16777216};
|
||||
uint numReads1 = 32;
|
||||
uint numReads2 = 256;
|
||||
uint sharedMemSizeBytes1 = sharedMemSize1 * sizeof(float);
|
||||
uint sharedMemSizeBytes2 = sharedMemSize2 * sizeof(float);
|
||||
int nIter = 1000;
|
||||
const unsigned threadsPerBlock = 64;
|
||||
|
||||
static int device = 0;
|
||||
HIP_CHECK(hipSetDevice(device));
|
||||
hipDeviceProp_t props;
|
||||
HIP_CHECK(hipGetDeviceProperties(&props, device));
|
||||
INFO("info: running on bus " << "0x" << props.pciBusID << " " << props.name
|
||||
<< " with " << props.multiProcessorCount << " CUs \n");
|
||||
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
for (int nTest = 0; nTest < numSizes; nTest++) {
|
||||
uint nBytes = Sizes[nTest % numSizes];
|
||||
ulong N = nBytes / sizeof(float);
|
||||
const unsigned blocks = N / threadsPerBlock;
|
||||
|
||||
hDst = new float[nBytes];
|
||||
HIP_CHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
memset(hDst, 0, nBytes);
|
||||
|
||||
HIP_CHECK(hipMalloc(&dDst, nBytes));
|
||||
HIP_CHECK(hipMemcpy(dDst, hDst, nBytes, hipMemcpyHostToDevice));
|
||||
|
||||
hipLaunchKernelGGL(sharedMemReadSpeed1, dim3(blocks),
|
||||
dim3(threadsPerBlock), 0, stream, dDst, N);
|
||||
HIP_CHECK(hipMemcpy(hDst, dDst, nBytes, hipMemcpyDeviceToHost));
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
int tmp = 0;
|
||||
for (int i = 0; i < N; i++) {
|
||||
if (i % threadsPerBlock == 0) {
|
||||
tmp = 0;
|
||||
}
|
||||
if (hDst[i] != tmp) {
|
||||
INFO("info: Data validation failed for warm up run! \n");
|
||||
INFO("info: expected " << tmp << " got " << hDst[i] << " \n");
|
||||
return false;
|
||||
}
|
||||
tmp += threadsPerBlock / 2;
|
||||
}
|
||||
|
||||
auto all_start = std::chrono::steady_clock::now();
|
||||
for (int i = 0; i < nIter; i++) {
|
||||
hipLaunchKernelGGL(sharedMemReadSpeed1, dim3(blocks),
|
||||
dim3(threadsPerBlock), 0, stream, dDst, N);
|
||||
}
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
auto all_end = std::chrono::steady_clock::now();
|
||||
std::chrono::duration<double> all_kernel_time = all_end - all_start;
|
||||
|
||||
// read speed in GB/s
|
||||
double perf = (static_cast<double>(blocks * threadsPerBlock)
|
||||
* (numReads1 * sizeof(float) + sharedMemSizeBytes1 / 64)
|
||||
* nIter * (1e-09)) / all_kernel_time.count();
|
||||
|
||||
INFO("info: read speed = " << std::setw(8) << perf << " GB/s for " <<
|
||||
sharedMemSizeBytes1 / 1024 << " KB shared memory with " <<
|
||||
std::setw(8) << blocks * threadsPerBlock << " threads, "
|
||||
<< std::setw(4) << numReads1 <<
|
||||
" reads in sharedMemReadSpeed1 kernel \n");
|
||||
|
||||
delete[] hDst;
|
||||
HIP_CHECK(hipFree(dDst));
|
||||
}
|
||||
|
||||
for (int nTest = 0; nTest < numSizes; nTest++) {
|
||||
uint nBytes = Sizes[nTest % numSizes];
|
||||
ulong N = nBytes / sizeof(float);
|
||||
const unsigned blocks = N / threadsPerBlock;
|
||||
|
||||
hDst = new float[nBytes];
|
||||
HIP_CHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
memset(hDst, 0, nBytes);
|
||||
|
||||
HIP_CHECK(hipMalloc(&dDst, nBytes));
|
||||
HIP_CHECK(hipMemcpy(dDst, hDst, nBytes, hipMemcpyHostToDevice));
|
||||
|
||||
hipLaunchKernelGGL(sharedMemReadSpeed2, dim3(blocks),
|
||||
dim3(threadsPerBlock), 0, stream, dDst, N);
|
||||
HIP_CHECK(hipMemcpy(hDst, dDst, nBytes, hipMemcpyDeviceToHost));
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
auto all_start = std::chrono::steady_clock::now();
|
||||
for (int i = 0; i < nIter; i++) {
|
||||
hipLaunchKernelGGL(sharedMemReadSpeed2, dim3(blocks),
|
||||
dim3(threadsPerBlock), 0, stream, dDst, N);
|
||||
}
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
auto all_end = std::chrono::steady_clock::now();
|
||||
std::chrono::duration<double> all_kernel_time = all_end - all_start;
|
||||
|
||||
// read speed in GB/s
|
||||
double perf = (static_cast<double>(blocks * threadsPerBlock)
|
||||
* (numReads2 * sizeof(float) + sharedMemSizeBytes2 / 64)
|
||||
* nIter * (1e-09)) / all_kernel_time.count();
|
||||
|
||||
INFO("info: read speed = " << std::setw(8) << perf << " GB/s for "
|
||||
<< sharedMemSizeBytes2 / 1024 << " KB shared memory with "
|
||||
<< std::setw(8) << blocks * threadsPerBlock << " threads, "
|
||||
<< std::setw(4) << numReads2 <<
|
||||
" reads in sharedMemReadSpeed2 kernel \n");
|
||||
|
||||
delete[] hDst;
|
||||
HIP_CHECK(hipFree(dDst));
|
||||
}
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
return true;
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Verify hipPerfSharedMemReadSpeed status.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - perftests/memory/hipPerfSharedMemReadSpeed.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.6
|
||||
*/
|
||||
|
||||
TEST_CASE("Perf_hipPerfSharedMemReadSpeed_test") {
|
||||
int numDevices = 0;
|
||||
HIP_CHECK(hipGetDeviceCount(&numDevices));
|
||||
|
||||
if (numDevices <= 0) {
|
||||
SUCCEED("Skipped testcase hipPerfSharedMemReadSpeed as"
|
||||
"there is no device to test.\n");
|
||||
} else {
|
||||
REQUIRE(true == hipPerfSharedMemReadSpeed_test());
|
||||
}
|
||||
}
|
||||
新しいイシューから参照
ユーザーをブロックする