diff --git a/projects/hip-tests/catch/include/resource_guards.hh b/projects/hip-tests/catch/include/resource_guards.hh index a98c828ef3..91d8b6ffd5 100644 --- a/projects/hip-tests/catch/include/resource_guards.hh +++ b/projects/hip-tests/catch/include/resource_guards.hh @@ -165,12 +165,18 @@ template class LinearAllocGuardMultiDim { const hipExtent extent_; }; -template class LinearAllocGuard2D : public LinearAllocGuardMultiDim { +template +class LinearAllocGuard2D : public LinearAllocGuardMultiDim { public: LinearAllocGuard2D(const size_t width_logical, const size_t height) : LinearAllocGuardMultiDim{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; diff --git a/projects/hip-tests/catch/unit/memory/hipDrvMemcpy2DUnaligned.cc b/projects/hip-tests/catch/unit/memory/hipDrvMemcpy2DUnaligned.cc index b5d2fb70f3..15e8cd1724 100644 --- a/projects/hip-tests/catch/unit/memory/hipDrvMemcpy2DUnaligned.cc +++ b/projects/hip-tests/catch/unit/memory/hipDrvMemcpy2DUnaligned.cc @@ -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 #include + +#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(DrvMemcpy2DUnalignedAdapter()); + } + SECTION("Peer access enabled") { + Memcpy2DDeviceToDeviceShell(DrvMemcpy2DUnalignedAdapter()); + } + } + + SECTION("Host to Device") { + Memcpy2DHostToDeviceShell(DrvMemcpy2DUnalignedAdapter()); + } + + SECTION("Device to Host") { + Memcpy2DDeviceToHostShell(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(DrvMemcpy2DUnalignedAdapter(), true); } + + SECTION("Device to Pinned Host") { + Memcpy2DDtoHPinnedSyncBehavior(DrvMemcpy2DUnalignedAdapter(), true); + } + + SECTION("Device to Pageable Host") { + Memcpy2DDtoHPageableSyncBehavior(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(DrvMemcpy2DUnalignedAdapter()); +} diff --git a/projects/hip-tests/catch/unit/memory/memcpy2d_tests_common.hh b/projects/hip-tests/catch/unit/memory/memcpy2d_tests_common.hh index 62d7667db6..54564dd9dd 100644 --- a/projects/hip-tests/catch/unit/memory/memcpy2d_tests_common.hh +++ b/projects/hip-tests/catch/unit/memory/memcpy2d_tests_common.hh @@ -30,14 +30,14 @@ THE SOFTWARE. #include #include -template +template 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 device_alloc(cols, rows); + LinearAllocGuard2D device_alloc(cols, rows); const size_t host_pitch = GENERATE_REF(device_alloc.width(), device_alloc.width() + 64); LinearAllocGuard 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 +template 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 src_alloc(cols * src_cols_mult, rows); + LinearAllocGuard2D src_alloc(cols * src_cols_mult, rows); HIP_CHECK(hipSetDevice(dst_device)); - LinearAllocGuard2D dst_alloc(cols, rows); + LinearAllocGuard2D dst_alloc(cols, rows); HIP_CHECK(hipSetDevice(src_device)); LinearAllocGuard 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 +template 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 device_alloc(cols, rows); + LinearAllocGuard2D 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 +template 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 host_alloc(host_alloc_type, 32 * sizeof(int) * 32); - LinearAllocGuard2D device_alloc(32, 32); + LinearAllocGuard2D 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 +template void Memcpy2DDtoHPageableSyncBehavior(F memcpy_func, const bool should_sync, const hipStream_t kernel_stream = nullptr) { LinearAllocGuard host_alloc(LinearAllocs::malloc, 32 * sizeof(int) * 32); - LinearAllocGuard2D device_alloc(32, 32); + LinearAllocGuard2D 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 +template void Memcpy2DDtoHPinnedSyncBehavior(F memcpy_func, const bool should_sync, const hipStream_t kernel_stream = nullptr) { LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, 32 * sizeof(int) * 32); - LinearAllocGuard2D device_alloc(32, 32); + LinearAllocGuard2D 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 +template void Memcpy2DDtoDSyncBehavior(F memcpy_func, const bool should_sync, const hipStream_t kernel_stream = nullptr) { - LinearAllocGuard2D src_alloc(32, 32); - LinearAllocGuard2D dst_alloc(32, 32); + LinearAllocGuard2D src_alloc(32, 32); + LinearAllocGuard2D 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 +template 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 device_alloc(cols, rows); + LinearAllocGuard2D device_alloc(cols, rows); LinearAllocGuard 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 src_alloc(cols, rows); - LinearAllocGuard2D dst_alloc(cols, rows); + LinearAllocGuard2D src_alloc(cols, rows); + LinearAllocGuard2D dst_alloc(cols, rows); LinearAllocGuard 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 device_alloc(cols, rows); + LinearAllocGuard2D device_alloc(cols, rows); LinearAllocGuard src_host_alloc(LinearAllocs::hipHostMalloc, device_alloc.width() * rows); LinearAllocGuard dst_host_alloc(LinearAllocs::hipHostMalloc, @@ -348,6 +348,70 @@ constexpr auto MemTypeUnified() { using PtrVariant = std::variant; +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(dst)) { + parms->dstMemoryType = MemTypeArray(); + parms->dstArray = std::get(dst); + } else { + parms->dstPitch = dpitch; + auto ptr = std::get(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(ptr); + break; + case hipMemcpyDefault: + parms->dstMemoryType = MemTypeUnified(); + parms->dstDevice = reinterpret_cast(ptr); + break; + default: + assert(false); + } + } + + if (std::holds_alternative(src)) { + parms->srcMemoryType = MemTypeArray(); + parms->srcArray = std::get(src); + } else { + parms->srcPitch = spitch; + auto ptr = std::get(src); + switch (kind) { + case hipMemcpyDeviceToHost: + case hipMemcpyDeviceToDevice: + parms->srcMemoryType = MemTypeDevice(); + parms->srcDevice = reinterpret_cast(ptr); + break; + case hipMemcpyHostToDevice: + case hipMemcpyHostToHost: + parms->srcMemoryType = MemTypeHost(); + parms->srcHost = ptr; + break; + case hipMemcpyDefault: + parms->srcMemoryType = MemTypeUnified(); + parms->srcDevice = reinterpret_cast(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 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(dst)) { - parms.dstMemoryType = MemTypeArray(); - parms.dstArray = std::get(dst); - } else { - parms.dstPitch = dpitch; - auto ptr = std::get(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(ptr); - break; - case hipMemcpyDefault: - parms.dstMemoryType = MemTypeUnified(); - parms.dstDevice = reinterpret_cast(ptr); - break; - default: - assert(false); - } - } - - if (std::holds_alternative(src)) { - parms.srcMemoryType = MemTypeArray(); - parms.srcArray = std::get(src); - } else { - parms.srcPitch = spitch; - auto ptr = std::get(src); - switch (kind) { - case hipMemcpyDeviceToHost: - case hipMemcpyDeviceToDevice: - parms.srcMemoryType = MemTypeDevice(); - parms.srcDevice = reinterpret_cast(ptr); - break; - case hipMemcpyHostToDevice: - case hipMemcpyHostToHost: - parms.srcMemoryType = MemTypeHost(); - parms.srcHost = ptr; - break; - case hipMemcpyDefault: - parms.srcMemoryType = MemTypeUnified(); - parms.srcDevice = reinterpret_cast(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 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); -} \ No newline at end of file +}