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

hipamd/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,