From 8a5f3ecb50db870cfbca1b2607127a6f004b90f0 Mon Sep 17 00:00:00 2001 From: foreman Date: Mon, 20 Jan 2020 19:05:16 -0500 Subject: [PATCH] P4 to Git Change 2060446 by vsytchen@vsytchen-ocl-win10-2 on 2020/01/20 18:58:37 SWDEV-215533 - [HIP][Windows]Output mismatch with tex3D() When calculating the pitch for hipMallocPitch() align it to the device's image pitch alignment. If we don't, PAL backend will overalign the pitch when creating an image buffer. ReviewBoardURL = http://ocltc.amd.com/reviews/r/18496/diff/ Affected files ... ... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#96 edit [ROCm/hip commit: d18baabc7078e82852af53d7071cd2717a5cecb6] --- projects/hip/api/hip/hip_memory.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/api/hip/hip_memory.cpp b/projects/hip/api/hip/hip_memory.cpp index 6b23913386..9cd5034eeb 100644 --- a/projects/hip/api/hip/hip_memory.cpp +++ b/projects/hip/api/hip/hip_memory.cpp @@ -340,7 +340,7 @@ hipError_t ihipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t heigh const amd::Image::Format imageFormat(*image_format); - *pitch = width * imageFormat.getElementSize(); + *pitch = amd::alignUp(width * imageFormat.getElementSize(), device->info().imagePitchAlignment_); size_t sizeBytes = *pitch * height * depth; *ptr = amd::SvmBuffer::malloc(*hip::getCurrentContext(), 0, sizeBytes,