AI/rocm-systems
2
0
Bifurcation 0
Code Tickets Demandes d'ajout Actions Paquets Projets Publications Wiki Activité

SWDEV-294420 - Ignore Image blit kernels if image instructions are not supported.

Parcourir la source 
Change-Id: I145172672b0b032aa722649b0c4ca9267e3e5c85
Cette révision appartient à :
kjayapra-amd
2021-09-28 12:34:24 -04:00
révisé par Karthik Jayaprakash
Parent d17108e8d0
révision 7413b7f79b
4 fichiers modifiés avec 149 ajouts et 98 suppressions
Télécharger le Fichier Patch Télécharger le Fichier des Différences Développer tous les fichiers Réduire tous les fichiers
rocclr/device/blitcl.cpp
+94 -77
Voir le fichier
@@ -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
rocclr/device/device.cpp
+8 -2
Voir le fichier
@@ -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<amd::Device*> 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;
rocclr/device/rocm/rocblit.cpp
+26 -2
Voir le fichier
@@ -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;
rocclr/device/rocm/rocblit.hpp
+21 -17
Voir le fichier
@@ -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,