AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

SWDEV-441603 - Add hipDrvMemcpy2DUnaligned test

Просмотр исходного кода 
Change-Id: I44b65867cabeb369fda8a82984fe3f694499942c


[ROCm/hip-tests commit: 639a9c2996]
Этот коммит содержится в:
Marko Arandjelovic
2024-03-19 13:28:45 +01:00
родитель 13c5e7a3e4
Коммит 48ee9f8ae7
3 изменённых файлов: 190 добавлений и 83 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
projects/hip-tests/catch/include/resource_guards.hh
+9 -3
Просмотреть файл
@@ -165,12 +165,18 @@ template <typename T> class LinearAllocGuardMultiDim {
const hipExtent extent_;
};
template <typename T> class LinearAllocGuard2D : public LinearAllocGuardMultiDim<T> {
template <typename T, bool unaligned = false>
class LinearAllocGuard2D : public LinearAllocGuardMultiDim<T> {
public:
LinearAllocGuard2D(const size_t width_logical, const size_t height)
: LinearAllocGuardMultiDim<T>{make_hipExtent(width_logical * sizeof(T), height, 1)} {
HIP_CHECK(hipMallocPitch(&this->pitched_ptr_.ptr, &this->pitched_ptr_.pitch,
this->extent_.width, this->extent_.height));
if (unaligned) {
this->pitched_ptr_.pitch = width_logical * sizeof(T);
HIP_CHECK(hipMalloc(&this->pitched_ptr_.ptr, this->pitched_ptr_.pitch * height));
} else {
HIP_CHECK(hipMallocPitch(&this->pitched_ptr_.ptr, &this->pitched_ptr_.pitch,
this->extent_.width, this->extent_.height));
}
}
LinearAllocGuard2D(const LinearAllocGuard2D&) = delete;
projects/hip-tests/catch/unit/memory/hipDrvMemcpy2DUnaligned.cc
+79
Просмотреть файл
@@ -16,8 +16,12 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip/hip_runtime_api.h>
#include <hip_test_common.hh>
#include "memcpy2d_tests_common.hh"
/**
* @addtogroup hipDrvMemcpy2DUnaligned hipDrvMemcpy2DUnaligned
* @{
@@ -252,3 +256,78 @@ TEST_CASE("Unit_hipDrvMemcpy2DUnaligned_FuncTst") {
free(dstH);
}
}
/**
* Test Description
* ------------------------
* - Basic test that copies and verifies copied data
* Test source
* ------------------------
* - unit/memory/hipDrvMemcpy2DUnaligned.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.0
*/
TEST_CASE("Unit_hipDrvMemcpy2DUnaligned_Positive_Basic") {
CHECK_IMAGE_SUPPORT
SECTION("Device to Device") {
SECTION("Peer access disabled") {
Memcpy2DDeviceToDeviceShell<false, false, true>(DrvMemcpy2DUnalignedAdapter());
}
SECTION("Peer access enabled") {
Memcpy2DDeviceToDeviceShell<false, true, true>(DrvMemcpy2DUnalignedAdapter());
}
}
SECTION("Host to Device") {
Memcpy2DHostToDeviceShell<false, true>(DrvMemcpy2DUnalignedAdapter());
}
SECTION("Device to Host") {
Memcpy2DDeviceToHostShell<false, true>(DrvMemcpy2DUnalignedAdapter());
}
}
/**
* Test Description
* ------------------------
* - Basic test that verifies synchronization behaviour
* Test source
* ------------------------
* - unit/memory/hipDrvMemcpy2DUnaligned.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.0
*/
TEST_CASE("Unit_hipDrvMemcpy2DUnaligned_Positive_Synchronization_Behavior") {
CHECK_IMAGE_SUPPORT
HIP_CHECK(hipDeviceSynchronize());
SECTION("Host to Device") { Memcpy2DHtoDSyncBehavior<true>(DrvMemcpy2DUnalignedAdapter(), true); }
SECTION("Device to Pinned Host") {
Memcpy2DDtoHPinnedSyncBehavior<true>(DrvMemcpy2DUnalignedAdapter(), true);
}
SECTION("Device to Pageable Host") {
Memcpy2DDtoHPageableSyncBehavior<true>(DrvMemcpy2DUnalignedAdapter(), true);
}
}
/**
* Test Description
* ------------------------
* - Basic test that copies and verifies copied data with zero width or height.
* Test source
* ------------------------
* - unit/memory/hipDrvMemcpy2DUnaligned.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.0
*/
TEST_CASE("Unit_hipDrvMemcpy2DUnaligned_Positive_Parameters") {
Memcpy2DZeroWidthHeight<false, true>(DrvMemcpy2DUnalignedAdapter());
}
projects/hip-tests/catch/unit/memory/memcpy2d_tests_common.hh
+102 -80
Просмотреть файл
@@ -30,14 +30,14 @@ THE SOFTWARE.
#include <resource_guards.hh>
#include <hip/driver_types.h>
template <bool should_synchronize, typename F>
template <bool should_synchronize, bool unaligned = false, typename F>
void Memcpy2DDeviceToHostShell(F memcpy_func, const hipStream_t kernel_stream = nullptr) {
const auto kind = GENERATE(hipMemcpyDeviceToHost, hipMemcpyDefault);
constexpr size_t cols = 127;
constexpr size_t rows = 128;
LinearAllocGuard2D<int> device_alloc(cols, rows);
LinearAllocGuard2D<int, unaligned> device_alloc(cols, rows);
const size_t host_pitch = GENERATE_REF(device_alloc.width(), device_alloc.width() + 64);
LinearAllocGuard<int> host_alloc(LinearAllocs::hipHostMalloc, host_pitch * rows);
@@ -59,7 +59,7 @@ void Memcpy2DDeviceToHostShell(F memcpy_func, const hipStream_t kernel_stream =
device_alloc.height(), 1, f);
}
template <bool should_synchronize, bool enable_peer_access, typename F>
template <bool should_synchronize, bool enable_peer_access, bool unaligned = false, typename F>
void Memcpy2DDeviceToDeviceShell(F memcpy_func, const hipStream_t kernel_stream = nullptr) {
const auto kind = GENERATE(hipMemcpyDeviceToDevice, hipMemcpyDefault);
@@ -89,9 +89,9 @@ void Memcpy2DDeviceToDeviceShell(F memcpy_func, const hipStream_t kernel_stream
HIP_CHECK(hipDeviceEnablePeerAccess(dst_device, 0));
}
LinearAllocGuard2D<int> src_alloc(cols * src_cols_mult, rows);
LinearAllocGuard2D<int, unaligned> src_alloc(cols * src_cols_mult, rows);
HIP_CHECK(hipSetDevice(dst_device));
LinearAllocGuard2D<int> dst_alloc(cols, rows);
LinearAllocGuard2D<int, unaligned> dst_alloc(cols, rows);
HIP_CHECK(hipSetDevice(src_device));
LinearAllocGuard<int> host_alloc(LinearAllocs::hipHostMalloc, dst_alloc.width() * rows);
@@ -116,14 +116,14 @@ void Memcpy2DDeviceToDeviceShell(F memcpy_func, const hipStream_t kernel_stream
dst_alloc.height(), 1, f);
}
template <bool should_synchronize, typename F>
template <bool should_synchronize, bool unaligned = false, typename F>
void Memcpy2DHostToDeviceShell(F memcpy_func, const hipStream_t kernel_stream = nullptr) {
const auto kind = GENERATE(hipMemcpyHostToDevice, hipMemcpyDefault);
constexpr size_t cols = 127;
constexpr size_t rows = 128;
LinearAllocGuard2D<int> device_alloc(cols, rows);
LinearAllocGuard2D<int, unaligned> device_alloc(cols, rows);
const size_t host_pitch = GENERATE_REF(device_alloc.pitch(), 2 * device_alloc.pitch());
@@ -188,46 +188,46 @@ void MemcpySyncBehaviorCheck(F memcpy_func, const bool should_sync,
}
}
template <typename F>
template <bool unaligned = false, typename F>
void Memcpy2DHtoDSyncBehavior(F memcpy_func, const bool should_sync,
const hipStream_t kernel_stream = nullptr) {
using LA = LinearAllocs;
const auto host_alloc_type = GENERATE(LA::malloc, LA::hipHostMalloc);
LinearAllocGuard<int> host_alloc(host_alloc_type, 32 * sizeof(int) * 32);
LinearAllocGuard2D<int> device_alloc(32, 32);
LinearAllocGuard2D<int, unaligned> device_alloc(32, 32);
MemcpySyncBehaviorCheck(std::bind(memcpy_func, device_alloc.ptr(), device_alloc.pitch(),
host_alloc.ptr(), device_alloc.width(), device_alloc.width(),
device_alloc.height(), hipMemcpyHostToDevice),
should_sync, kernel_stream);
}
template <typename F>
template <bool unaligned = false, typename F>
void Memcpy2DDtoHPageableSyncBehavior(F memcpy_func, const bool should_sync,
const hipStream_t kernel_stream = nullptr) {
LinearAllocGuard<int> host_alloc(LinearAllocs::malloc, 32 * sizeof(int) * 32);
LinearAllocGuard2D<int> device_alloc(32, 32);
LinearAllocGuard2D<int, unaligned> device_alloc(32, 32);
MemcpySyncBehaviorCheck(std::bind(memcpy_func, host_alloc.ptr(), device_alloc.width(),
device_alloc.ptr(), device_alloc.pitch(), device_alloc.width(),
device_alloc.height(), hipMemcpyDeviceToHost),
should_sync, kernel_stream);
}
template <typename F>
template <bool unaligned = false, typename F>
void Memcpy2DDtoHPinnedSyncBehavior(F memcpy_func, const bool should_sync,
const hipStream_t kernel_stream = nullptr) {
LinearAllocGuard<int> host_alloc(LinearAllocs::hipHostMalloc, 32 * sizeof(int) * 32);
LinearAllocGuard2D<int> device_alloc(32, 32);
LinearAllocGuard2D<int, unaligned> device_alloc(32, 32);
MemcpySyncBehaviorCheck(std::bind(memcpy_func, host_alloc.ptr(), device_alloc.width(),
device_alloc.ptr(), device_alloc.pitch(), device_alloc.width(),
device_alloc.height(), hipMemcpyDeviceToHost),
should_sync, kernel_stream);
}
template <typename F>
template <bool unaligned = false, typename F>
void Memcpy2DDtoDSyncBehavior(F memcpy_func, const bool should_sync,
const hipStream_t kernel_stream = nullptr) {
LinearAllocGuard2D<int> src_alloc(32, 32);
LinearAllocGuard2D<int> dst_alloc(32, 32);
LinearAllocGuard2D<int, unaligned> src_alloc(32, 32);
LinearAllocGuard2D<int, unaligned> dst_alloc(32, 32);
MemcpySyncBehaviorCheck(
std::bind(memcpy_func, dst_alloc.ptr(), dst_alloc.pitch(), src_alloc.ptr(), src_alloc.pitch(),
dst_alloc.width(), dst_alloc.height(), hipMemcpyDeviceToDevice),
@@ -248,7 +248,7 @@ void Memcpy2DHtoHSyncBehavior(F memcpy_func, const bool should_sync,
should_sync, kernel_stream);
}
template <bool should_synchronize, typename F>
template <bool should_synchronize, bool unaligned = false, typename F>
void Memcpy2DZeroWidthHeight(F memcpy_func, const hipStream_t stream = nullptr) {
constexpr size_t cols = 63;
constexpr size_t rows = 64;
@@ -257,7 +257,7 @@ void Memcpy2DZeroWidthHeight(F memcpy_func, const hipStream_t stream = nullptr)
GENERATE(std::make_pair(0, 1), std::make_pair(1, 0), std::make_pair(0, 0));
SECTION("Device to Host") {
LinearAllocGuard2D<uint8_t> device_alloc(cols, rows);
LinearAllocGuard2D<uint8_t, unaligned> device_alloc(cols, rows);
LinearAllocGuard<uint8_t> host_alloc(LinearAllocs::hipHostMalloc, device_alloc.width() * rows);
std::fill_n(host_alloc.ptr(), device_alloc.width_logical() * device_alloc.height(), 42);
HIP_CHECK(hipMemset2D(device_alloc.ptr(), device_alloc.pitch(), 1, device_alloc.width(),
@@ -274,8 +274,8 @@ void Memcpy2DZeroWidthHeight(F memcpy_func, const hipStream_t stream = nullptr)
}
SECTION("Device to Device") {
LinearAllocGuard2D<uint8_t> src_alloc(cols, rows);
LinearAllocGuard2D<uint8_t> dst_alloc(cols, rows);
LinearAllocGuard2D<uint8_t, unaligned> src_alloc(cols, rows);
LinearAllocGuard2D<uint8_t, unaligned> dst_alloc(cols, rows);
LinearAllocGuard<uint8_t> host_alloc(LinearAllocs::hipHostMalloc, dst_alloc.width() * rows);
HIP_CHECK(
hipMemset2D(src_alloc.ptr(), src_alloc.pitch(), 1, src_alloc.width(), src_alloc.height()));
@@ -294,7 +294,7 @@ void Memcpy2DZeroWidthHeight(F memcpy_func, const hipStream_t stream = nullptr)
}
SECTION("Host to Device") {
LinearAllocGuard2D<uint8_t> device_alloc(cols, rows);
LinearAllocGuard2D<uint8_t, unaligned> device_alloc(cols, rows);
LinearAllocGuard<uint8_t> src_host_alloc(LinearAllocs::hipHostMalloc,
device_alloc.width() * rows);
LinearAllocGuard<uint8_t> dst_host_alloc(LinearAllocs::hipHostMalloc,
@@ -348,6 +348,70 @@ constexpr auto MemTypeUnified() {
using PtrVariant = std::variant<void*, hipArray_t>;
constexpr void InitializeMemcpy2DParams(hip_Memcpy2D* parms, PtrVariant dst, size_t dpitch,
PtrVariant src, size_t spitch, size_t width, size_t height,
hipMemcpyKind kind, hipExtent src_offset = {0, 0, 0},
hipExtent dst_offset = {0, 0, 0}) {
if (std::holds_alternative<hipArray_t>(dst)) {
parms->dstMemoryType = MemTypeArray();
parms->dstArray = std::get<hipArray_t>(dst);
} else {
parms->dstPitch = dpitch;
auto ptr = std::get<void*>(dst);
switch (kind) {
case hipMemcpyDeviceToHost:
case hipMemcpyHostToHost:
parms->dstMemoryType = MemTypeHost();
parms->dstHost = ptr;
break;
case hipMemcpyDeviceToDevice:
case hipMemcpyHostToDevice:
parms->dstMemoryType = MemTypeDevice();
parms->dstDevice = reinterpret_cast<hipDeviceptr_t>(ptr);
break;
case hipMemcpyDefault:
parms->dstMemoryType = MemTypeUnified();
parms->dstDevice = reinterpret_cast<hipDeviceptr_t>(ptr);
break;
default:
assert(false);
}
}
if (std::holds_alternative<hipArray_t>(src)) {
parms->srcMemoryType = MemTypeArray();
parms->srcArray = std::get<hipArray_t>(src);
} else {
parms->srcPitch = spitch;
auto ptr = std::get<void*>(src);
switch (kind) {
case hipMemcpyDeviceToHost:
case hipMemcpyDeviceToDevice:
parms->srcMemoryType = MemTypeDevice();
parms->srcDevice = reinterpret_cast<hipDeviceptr_t>(ptr);
break;
case hipMemcpyHostToDevice:
case hipMemcpyHostToHost:
parms->srcMemoryType = MemTypeHost();
parms->srcHost = ptr;
break;
case hipMemcpyDefault:
parms->srcMemoryType = MemTypeUnified();
parms->srcDevice = reinterpret_cast<hipDeviceptr_t>(ptr);
break;
default:
assert(false);
}
}
parms->WidthInBytes = width;
parms->Height = height;
parms->srcXInBytes = src_offset.width;
parms->srcY = src_offset.height;
parms->dstXInBytes = dst_offset.width;
parms->dstY = dst_offset.height;
}
template <bool async = false>
constexpr auto MemcpyParam2DAdapter(const hipExtent src_offset = {0, 0, 0},
const hipExtent dst_offset = {0, 0, 0}) {
@@ -356,64 +420,8 @@ constexpr auto MemcpyParam2DAdapter(const hipExtent src_offset = {0, 0, 0},
hip_Memcpy2D parms = {};
memset(&parms, 0x0, sizeof(hip_Memcpy2D));
if (std::holds_alternative<hipArray_t>(dst)) {
parms.dstMemoryType = MemTypeArray();
parms.dstArray = std::get<hipArray_t>(dst);
} else {
parms.dstPitch = dpitch;
auto ptr = std::get<void*>(dst);
switch (kind) {
case hipMemcpyDeviceToHost:
case hipMemcpyHostToHost:
parms.dstMemoryType = MemTypeHost();
parms.dstHost = ptr;
break;
case hipMemcpyDeviceToDevice:
case hipMemcpyHostToDevice:
parms.dstMemoryType = MemTypeDevice();
parms.dstDevice = reinterpret_cast<hipDeviceptr_t>(ptr);
break;
case hipMemcpyDefault:
parms.dstMemoryType = MemTypeUnified();
parms.dstDevice = reinterpret_cast<hipDeviceptr_t>(ptr);
break;
default:
assert(false);
}
}
if (std::holds_alternative<hipArray_t>(src)) {
parms.srcMemoryType = MemTypeArray();
parms.srcArray = std::get<hipArray_t>(src);
} else {
parms.srcPitch = spitch;
auto ptr = std::get<void*>(src);
switch (kind) {
case hipMemcpyDeviceToHost:
case hipMemcpyDeviceToDevice:
parms.srcMemoryType = MemTypeDevice();
parms.srcDevice = reinterpret_cast<hipDeviceptr_t>(ptr);
break;
case hipMemcpyHostToDevice:
case hipMemcpyHostToHost:
parms.srcMemoryType = MemTypeHost();
parms.srcHost = ptr;
break;
case hipMemcpyDefault:
parms.srcMemoryType = MemTypeUnified();
parms.srcDevice = reinterpret_cast<hipDeviceptr_t>(ptr);
break;
default:
assert(false);
}
}
parms.WidthInBytes = width;
parms.Height = height;
parms.srcXInBytes = src_offset.width;
parms.srcY = src_offset.height;
parms.dstXInBytes = dst_offset.width;
parms.dstY = dst_offset.height;
InitializeMemcpy2DParams(&parms, dst, dpitch, src, spitch, width, height, kind, src_offset,
dst_offset);
if constexpr (async) {
return hipMemcpyParam2DAsync(&parms, stream);
@@ -423,6 +431,20 @@ constexpr auto MemcpyParam2DAdapter(const hipExtent src_offset = {0, 0, 0},
};
}
#if HT_AMD
constexpr auto DrvMemcpy2DUnalignedAdapter() {
return [=](PtrVariant dst, size_t dpitch, PtrVariant src, size_t spitch, size_t width,
size_t height, hipMemcpyKind kind, hipStream_t stream = nullptr) {
(void)stream;
hip_Memcpy2D parms = {};
memset(&parms, 0x0, sizeof(hip_Memcpy2D));
InitializeMemcpy2DParams(&parms, dst, dpitch, src, spitch, width, height, kind);
return hipDrvMemcpy2DUnaligned(&parms);
};
}
#endif
template <bool should_synchronize, typename F>
void MemcpyParam2DArrayHostShell(F memcpy_func, const hipStream_t kernel_stream = nullptr) {
hipExtent extent{127 * sizeof(int), 128, 1};
@@ -521,4 +543,4 @@ void MemcpyParam2DArrayDeviceShell(F memcpy_func, const hipStream_t kernel_strea
};
PitchedMemoryVerify(host_alloc.ptr(), extent.width, extent.width / sizeof(int), extent.height,
extent.depth, f);
}
}