From 7413b7f79beb702a8c80de885ea058ebc8e1cedd Mon Sep 17 00:00:00 2001 From: kjayapra-amd Date: Tue, 28 Sep 2021 12:34:24 -0400 Subject: [PATCH] SWDEV-294420 - Ignore Image blit kernels if image instructions are not supported. Change-Id: I145172672b0b032aa722649b0c4ca9267e3e5c85 --- rocclr/device/blitcl.cpp | 171 ++++++++++++++++++--------------- rocclr/device/device.cpp | 10 +- rocclr/device/rocm/rocblit.cpp | 28 +++++- rocclr/device/rocm/rocblit.hpp | 38 ++++---- 4 files changed, 149 insertions(+), 98 deletions(-) diff --git a/rocclr/device/blitcl.cpp b/rocclr/device/blitcl.cpp index c5d5830bcd..13df6fcc0f 100644 --- a/rocclr/device/blitcl.cpp +++ b/rocclr/device/blitcl.cpp @@ -22,105 +22,122 @@ namespace device { #define BLIT_KERNELS(...) #__VA_ARGS__ -const char* BlitSourceCode = BLIT_KERNELS( - extern void __amd_copyBufferRect(__global uchar*, __global uchar*, ulong4, ulong4, ulong4); +const char* BlitLinearSourceCode = BLIT_KERNELS( + // Extern + extern void __amd_fillBufferAligned(__global uchar*, __global ushort*, __global uint*, + __global ulong*, __constant uchar*, uint, ulong, ulong); - extern void __amd_copyBufferRectAligned(__global uint*, __global uint*, ulong4, ulong4, ulong4); + extern void __amd_copyBuffer(__global uchar*, __global uchar*, ulong, ulong, ulong, uint); - extern void __amd_copyBuffer(__global uchar*, __global uchar*, ulong, ulong, ulong, uint); + extern void __amd_copyBufferAligned(__global uint*, __global uint*, ulong, ulong, ulong, uint); - extern void __amd_copyBufferAligned(__global uint*, __global uint*, ulong, ulong, ulong, uint); + extern void __amd_copyBufferRect(__global uchar*, __global uchar*, ulong4, ulong4, ulong4); - extern void __amd_fillBufferAligned(__global uchar*, __global ushort*, __global uint*, __global ulong*, - __constant uchar*, uint, ulong, ulong); + extern void __amd_copyBufferRectAligned(__global uint*, __global uint*, ulong4, ulong4, ulong4); - extern void __amd_streamOpsWrite(__global uint*, __global ulong*, ulong, ulong); + extern void __amd_streamOpsWrite(__global uint*, __global ulong*, ulong, ulong); - extern void __amd_streamOpsWait(__global uint*,__global ulong*, ulong, ulong, ulong); + extern void __amd_streamOpsWait(__global uint*,__global ulong*, ulong, ulong, ulong); - __kernel void __amd_rocclr_copyBufferRect(__global uchar* src, __global uchar* dst, ulong4 srcRect, - ulong4 dstRect, ulong4 size) { - __amd_copyBufferRect(src, dst, srcRect, dstRect, size); - } + // Implementation + __kernel void __amd_rocclr_fillBufferAligned(__global uchar* bufUChar, + __global ushort* bufUShort, + __global uint* bufUInt, + __global ulong* bufULong, + __constant uchar* pattern, + uint patternSize, ulong offset, + ulong size) { + __amd_fillBufferAligned(bufUChar, bufUShort, bufUInt, bufULong, + pattern, patternSize, offset, size); + } - __kernel void __amd_rocclr_copyBufferRectAligned(__global uint* src, __global uint* dst, ulong4 srcRect, - ulong4 dstRect, ulong4 size) { - __amd_copyBufferRectAligned(src, dst, srcRect, dstRect, size); - } + __kernel void __amd_rocclr_copyBuffer(__global uchar* srcI, __global uchar* dstI, + ulong srcOrigin, ulong dstOrigin, ulong size, + uint remain) { + __amd_copyBuffer(srcI, dstI, srcOrigin, dstOrigin, size, remain); + } - __kernel void __amd_rocclr_copyBuffer(__global uchar* srcI, __global uchar* dstI, ulong srcOrigin, - ulong dstOrigin, ulong size, uint remain) { - __amd_copyBuffer(srcI, dstI, srcOrigin, dstOrigin, size, remain); - } + __kernel void __amd_rocclr_copyBufferAligned(__global uint* src, __global uint* dst, + ulong srcOrigin, ulong dstOrigin, ulong size, + uint alignment) { + __amd_copyBufferAligned(src, dst, srcOrigin, dstOrigin, size, alignment); + } - __kernel void __amd_rocclr_copyBufferAligned(__global uint* src, __global uint* dst, ulong srcOrigin, - ulong dstOrigin, ulong size, uint alignment) { - __amd_copyBufferAligned(src, dst, srcOrigin, dstOrigin, size, alignment); - } + __kernel void __amd_rocclr_copyBufferRect(__global uchar* src, __global uchar* dst, + ulong4 srcRect, ulong4 dstRect, ulong4 size) { + __amd_copyBufferRect(src, dst, srcRect, dstRect, size); + } - __kernel void __amd_rocclr_fillBufferAligned(__global uchar* bufUChar, - __global ushort* bufUShort, - __global uint* bufUInt, - __global ulong* bufULong, - __constant uchar* pattern, - uint patternSize, ulong offset, - ulong size) { - __amd_fillBufferAligned(bufUChar, bufUShort, bufUInt, bufULong, - pattern, patternSize, offset, size); - } + __kernel void __amd_rocclr_copyBufferRectAligned(__global uint* src, __global uint* dst, + ulong4 srcRect, ulong4 dstRect, ulong4 size) { + __amd_copyBufferRectAligned(src, dst, srcRect, dstRect, size); + } - extern void __amd_copyBufferToImage(__global uint*, __write_only image2d_array_t, ulong4, - int4, int4, uint4, ulong4); + __kernel void __amd_rocclr_streamOpsWrite(__global uint* ptrInt, __global ulong* ptrUlong, + ulong value, ulong sizeBytes) { + __amd_streamOpsWrite(ptrInt, ptrUlong, value, sizeBytes); + } - extern void __amd_copyImageToBuffer(__read_only image2d_array_t, __global uint*, - __global ushort*, __global uchar*, int4, ulong4, int4, - uint4, ulong4); + __kernel void __amd_rocclr_streamOpsWait(__global uint* ptrInt, __global ulong* ptrUlong, + ulong value, ulong flags, ulong mask) { + __amd_streamOpsWait(ptrInt, ptrUlong, value, flags, mask); + } - extern void __amd_copyImage(__read_only image2d_array_t, __write_only image2d_array_t, int4, - int4, int4); +); - extern void __amd_copyImage1DA(__read_only image2d_array_t, __write_only image2d_array_t, int4, - int4, int4); +const char* BlitImageSourceCode = BLIT_KERNELS( + // Extern + extern void __amd_fillImage(__write_only image2d_array_t, float4, int4, uint4, int4, int4, + uint); - extern void __amd_fillImage(__write_only image2d_array_t, float4, int4, uint4, int4, int4, - uint); + extern void __amd_copyImage(__read_only image2d_array_t, __write_only image2d_array_t, int4, + int4, int4); + extern void __amd_copyImage1DA(__read_only image2d_array_t, __write_only image2d_array_t, int4, + int4, int4); - __kernel void __amd_rocclr_copyBufferToImage(__global uint* src, __write_only image2d_array_t dst, - ulong4 srcOrigin, int4 dstOrigin, int4 size, uint4 format, - ulong4 pitch) { - __amd_copyBufferToImage(src, dst, srcOrigin, dstOrigin, size, format, pitch); - } + extern void __amd_copyBufferToImage(__global uint*, __write_only image2d_array_t, ulong4, + int4, int4, uint4, ulong4); - __kernel void __amd_rocclr_copyImageToBuffer(__read_only image2d_array_t src, __global uint* dstUInt, - __global ushort* dstUShort, __global uchar* dstUChar, - int4 srcOrigin, ulong4 dstOrigin, int4 size, uint4 format, - ulong4 pitch) { - __amd_copyImageToBuffer(src, dstUInt, dstUShort, dstUChar, srcOrigin, dstOrigin, size, format, - pitch); - } + extern void __amd_copyImageToBuffer(__read_only image2d_array_t, __global uint*, + __global ushort*, __global uchar*, int4, ulong4, int4, + uint4, ulong4); - __kernel void __amd_rocclr_copyImage(__read_only image2d_array_t src, __write_only image2d_array_t dst, - int4 srcOrigin, int4 dstOrigin, - int4 size) { __amd_copyImage(src, dst, srcOrigin, dstOrigin, size); } + __kernel void __amd_rocclr_fillImage(__write_only image2d_array_t image, float4 patternFLOAT4, + int4 patternINT4, uint4 patternUINT4, int4 origin, + int4 size, uint type) { + __amd_fillImage(image, patternFLOAT4, patternINT4, patternUINT4, origin, size, type); + } - __kernel void __amd_rocclr_copyImage1DA(__read_only image2d_array_t src, __write_only image2d_array_t dst, - int4 srcOrigin, int4 dstOrigin, int4 size) { - __amd_copyImage1DA(src, dst, srcOrigin, dstOrigin, size); - } + __kernel void __amd_rocclr_copyImage(__read_only image2d_array_t src, + __write_only image2d_array_t dst, int4 srcOrigin, + int4 dstOrigin, int4 size) { + __amd_copyImage(src, dst, srcOrigin, dstOrigin, size); + } - __kernel void __amd_rocclr_fillImage(__write_only image2d_array_t image, float4 patternFLOAT4, - int4 patternINT4, uint4 patternUINT4, int4 origin, int4 size, - uint type) { - __amd_fillImage(image, patternFLOAT4, patternINT4, patternUINT4, origin, size, type); - } + __kernel void __amd_rocclr_copyImage1DA(__read_only image2d_array_t src, + __write_only image2d_array_t dst, int4 srcOrigin, + int4 dstOrigin, int4 size) { + __amd_copyImage1DA(src, dst, srcOrigin, dstOrigin, size); + } - __kernel void __amd_rocclr_streamOpsWait(__global uint* ptrInt, __global ulong* ptrUlong, ulong value, ulong flags, - ulong mask) { - __amd_streamOpsWait(ptrInt, ptrUlong, value, flags, mask); - } + __kernel void __amd_rocclr_copyBufferToImage(__global uint* src, + __write_only image2d_array_t dst, ulong4 srcOrigin, + int4 dstOrigin, int4 size, uint4 format, + ulong4 pitch) { + __amd_copyBufferToImage(src, dst, srcOrigin, dstOrigin, size, format, pitch); + } + + __kernel void __amd_rocclr_copyImageToBuffer(__read_only image2d_array_t src, + __global uint* dstUInt, + __global ushort* dstUShort, + __global uchar* dstUChar, + int4 srcOrigin, ulong4 dstOrigin, int4 size, + uint4 format, ulong4 pitch) { + __amd_copyImageToBuffer(src, dstUInt, dstUShort, dstUChar, srcOrigin, dstOrigin, size, format, + pitch); + } + +); - __kernel void __amd_rocclr_streamOpsWrite(__global uint* ptrInt, __global ulong* ptrUlong, ulong value, ulong sizeBytes) { - __amd_streamOpsWrite(ptrInt, ptrUlong, value, sizeBytes); - }); } // namespace device diff --git a/rocclr/device/device.cpp b/rocclr/device/device.cpp index b9723c45f8..57f6fb5d49 100644 --- a/rocclr/device/device.cpp +++ b/rocclr/device/device.cpp @@ -79,7 +79,8 @@ constexpr char hsaIsaNamePrefix[] = "amdgcn-amd-amdhsa--"; } // namespace namespace device { -extern const char* BlitSourceCode; +extern const char* BlitLinearSourceCode; +extern const char* BlitImageSourceCode; bool VirtualDevice::ActiveWait() const { return device_().ActiveWait(); @@ -359,7 +360,12 @@ bool Device::BlitProgram::create(amd::Device* device, const char* extraKernels, const char* extraOptions) { std::vector devices; devices.push_back(device); - std::string kernels(device::BlitSourceCode); + std::string kernels(device::BlitLinearSourceCode); + std::string image_kernels(device::BlitImageSourceCode); + + if (device->info().imageSupport_) { + kernels += image_kernels; + } if (extraKernels != nullptr) { kernels += extraKernels; diff --git a/rocclr/device/rocm/rocblit.cpp b/rocclr/device/rocm/rocblit.cpp index 868258689a..dab748249d 100644 --- a/rocclr/device/rocm/rocblit.cpp +++ b/rocclr/device/rocm/rocblit.cpp @@ -786,7 +786,7 @@ KernelBlitManager::KernelBlitManager(VirtualGPU& gpu, Setup setup) } KernelBlitManager::~KernelBlitManager() { - for (uint i = 0; i < BlitTotal; ++i) { + for (uint i = 0; i < NumBlitKernels(); ++i) { if (nullptr != kernels_[i]) { kernels_[i]->release(); } @@ -835,7 +835,7 @@ bool KernelBlitManager::createProgram(Device& device) { bool result = false; do { // Create kernel objects for all blits - for (uint i = 0; i < BlitTotal; ++i) { + for (uint i = 0; i < NumBlitKernels(); ++i) { const amd::Symbol* symbol = program_->findSymbol(BlitName[i]); if (symbol == nullptr) { break; @@ -896,6 +896,9 @@ bool KernelBlitManager::copyBufferToImage(device::Memory& srcMemory, device::Mem const amd::Coord3D& srcOrigin, const amd::Coord3D& dstOrigin, const amd::Coord3D& size, bool entire, size_t rowPitch, size_t slicePitch) const { + + guarantee((dev().info().imageSupport_ != false), "Image not supported on this device"); + amd::ScopedLock k(lockXferOps_); bool result = false; static const bool CopyRect = false; @@ -967,6 +970,9 @@ bool KernelBlitManager::copyBufferToImageKernel(device::Memory& srcMemory, const amd::Coord3D& dstOrigin, const amd::Coord3D& size, bool entire, size_t rowPitch, size_t slicePitch) const { + + guarantee((dev().info().imageSupport_ != false), "Image not supported on this device"); + bool rejected = false; Memory* dstView = &gpuMem(dstMemory); bool releaseView = false; @@ -1111,6 +1117,9 @@ bool KernelBlitManager::copyImageToBuffer(device::Memory& srcMemory, device::Mem const amd::Coord3D& srcOrigin, const amd::Coord3D& dstOrigin, const amd::Coord3D& size, bool entire, size_t rowPitch, size_t slicePitch) const { + + guarantee((dev().info().imageSupport_ != false), "Image not supported on this device"); + amd::ScopedLock k(lockXferOps_); bool result = false; static const bool CopyRect = false; @@ -1160,6 +1169,9 @@ bool KernelBlitManager::copyImageToBufferKernel(device::Memory& srcMemory, const amd::Coord3D& dstOrigin, const amd::Coord3D& size, bool entire, size_t rowPitch, size_t slicePitch) const { + + guarantee((dev().info().imageSupport_ != false), "Image not supported on this device"); + bool rejected = false; Memory* srcView = &gpuMem(srcMemory); bool releaseView = false; @@ -1308,6 +1320,9 @@ bool KernelBlitManager::copyImageToBufferKernel(device::Memory& srcMemory, bool KernelBlitManager::copyImage(device::Memory& srcMemory, device::Memory& dstMemory, const amd::Coord3D& srcOrigin, const amd::Coord3D& dstOrigin, const amd::Coord3D& size, bool entire) const { + + guarantee((dev().info().imageSupport_ != false), "Image not supported on this device"); + amd::ScopedLock k(lockXferOps_); bool rejected = false; Memory* srcView = &gpuMem(srcMemory); @@ -1469,6 +1484,9 @@ void FindPinSize(size_t& pinSize, const amd::Coord3D& size, size_t& rowPitch, si bool KernelBlitManager::readImage(device::Memory& srcMemory, void* dstHost, const amd::Coord3D& origin, const amd::Coord3D& size, size_t rowPitch, size_t slicePitch, bool entire) const { + + guarantee((dev().info().imageSupport_ != false), "Image not supported on this device"); + amd::ScopedLock k(lockXferOps_); bool result = false; @@ -1516,6 +1534,9 @@ bool KernelBlitManager::readImage(device::Memory& srcMemory, void* dstHost, bool KernelBlitManager::writeImage(const void* srcHost, device::Memory& dstMemory, const amd::Coord3D& origin, const amd::Coord3D& size, size_t rowPitch, size_t slicePitch, bool entire) const { + + guarantee((dev().info().imageSupport_ != false), "Image not supported on this device"); + amd::ScopedLock k(lockXferOps_); bool result = false; @@ -2109,6 +2130,9 @@ bool KernelBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& ds bool KernelBlitManager::fillImage(device::Memory& memory, const void* pattern, const amd::Coord3D& origin, const amd::Coord3D& size, bool entire) const { + + guarantee((dev().info().imageSupport_ != false), "Image not supported on this device"); + amd::ScopedLock k(lockXferOps_); bool result = false; diff --git a/rocclr/device/rocm/rocblit.hpp b/rocclr/device/rocm/rocblit.hpp index 3de0d75bc2..217bb8ee5d 100644 --- a/rocclr/device/rocm/rocblit.hpp +++ b/rocclr/device/rocm/rocblit.hpp @@ -230,20 +230,21 @@ class DmaBlitManager : public device::HostBlitManager { class KernelBlitManager : public DmaBlitManager { public: enum { - BlitCopyImage = 0, + FillBufferAligned = 0, + BlitCopyBuffer, + BlitCopyBufferAligned, + BlitCopyBufferRect, + BlitCopyBufferRectAligned, + StreamOpsWrite, + StreamOpsWait, + Scheduler, + GwsInit, + BlitLinearTotal, + FillImage = BlitLinearTotal, + BlitCopyImage, BlitCopyImage1DA, BlitCopyImageToBuffer, BlitCopyBufferToImage, - BlitCopyBufferRect, - BlitCopyBufferRectAligned, - BlitCopyBuffer, - BlitCopyBufferAligned, - FillBufferAligned, - FillImage, - Scheduler, - GwsInit, - StreamOpsWrite, - StreamOpsWait, BlitTotal }; @@ -462,6 +463,10 @@ class KernelBlitManager : public DmaBlitManager { return constantBufferOffset_; } + inline uint32_t NumBlitKernels() { + return (dev().info().imageSupport_) ? BlitTotal : BlitLinearTotal; + } + //! Disable copy constructor KernelBlitManager(const KernelBlitManager&); @@ -477,12 +482,11 @@ class KernelBlitManager : public DmaBlitManager { }; static const char* BlitName[KernelBlitManager::BlitTotal] = { - "__amd_rocclr_copyImage", "__amd_rocclr_copyImage1DA", "__amd_rocclr_copyImageToBuffer", - "__amd_rocclr_copyBufferToImage", "__amd_rocclr_copyBufferRect", - "__amd_rocclr_copyBufferRectAligned", "__amd_rocclr_copyBuffer", - "__amd_rocclr_copyBufferAligned", "__amd_rocclr_fillBufferAligned", - "__amd_rocclr_fillImage", "__amd_rocclr_scheduler", - "__amd_rocclr_gwsInit", "__amd_rocclr_streamOpsWrite", "__amd_rocclr_streamOpsWait" + "__amd_rocclr_fillBufferAligned", "__amd_rocclr_copyBuffer", "__amd_rocclr_copyBufferAligned", + "__amd_rocclr_copyBufferRect", "__amd_rocclr_copyBufferRectAligned", + "__amd_rocclr_streamOpsWrite", "__amd_rocclr_streamOpsWait", "__amd_rocclr_scheduler", + "__amd_rocclr_gwsInit", "__amd_rocclr_fillImage", "__amd_rocclr_copyImage", + "__amd_rocclr_copyImage1DA", "__amd_rocclr_copyImageToBuffer", "__amd_rocclr_copyBufferToImage", }; inline void KernelBlitManager::setArgument(amd::Kernel* kernel, size_t index,