foreman
824a34cf54
SWDEV-107442, SWDEV-108405 - Need to release GPU memory fence before releasing resources When we call submitKernelInternal, we need to release GPU memory fence before release resources used in the kernel otherwise we can page fault. ReviewBoardURL = http://ocltc.amd.com/reviews/r/11955/diff/ Affected files ... ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocblit.cpp#5 edit
1538 строки
52 KiB
C++
1538 строки
52 KiB
C++
//
|
|
// Copyright (c) 2013 Advanced Micro Devices, Inc. All rights reserved.
|
|
//
|
|
|
|
#include "platform/commandqueue.hpp"
|
|
#include "device/rocm/rocdevice.hpp"
|
|
#include "device/rocm/rocblit.hpp"
|
|
#include "device/rocm/rocmemory.hpp"
|
|
#include "device/rocm/rocvirtual.hpp"
|
|
#include "utils/debug.hpp"
|
|
|
|
namespace roc {
|
|
|
|
|
|
void
|
|
FindPinSize(
|
|
size_t& pinSize, const amd::Coord3D& size,
|
|
size_t& rowPitch, size_t& slicePitch, const Image& image)
|
|
{
|
|
size_t elementSize = image.owner()->asImage()->getImageFormat().getElementSize();
|
|
pinSize = size[0] * elementSize;
|
|
if ((rowPitch == 0) || (rowPitch == pinSize)) {
|
|
rowPitch = 0;
|
|
}
|
|
else {
|
|
pinSize = rowPitch;
|
|
}
|
|
|
|
// Calculate the pin size, which should be equal to the copy size
|
|
for (uint i = 1; i < 3; ++i) {
|
|
pinSize *= size[i];
|
|
if (i == 1) {
|
|
if ((slicePitch == 0) || (slicePitch == pinSize)) {
|
|
slicePitch = 0;
|
|
}
|
|
else {
|
|
if (image.getHsaImageDescriptor().geometry != HSA_EXT_IMAGE_GEOMETRY_1DA) {
|
|
pinSize = slicePitch;
|
|
}
|
|
else {
|
|
pinSize = slicePitch * size[i];
|
|
}
|
|
}
|
|
}
|
|
}
|
|
}
|
|
|
|
HsaBlitManager::HsaBlitManager(device::VirtualDevice& vDev, Setup setup)
|
|
: HostBlitManager(vDev, setup),
|
|
roc_device_(reinterpret_cast<const roc::Device &>(dev_)) {
|
|
completion_signal_.handle = 0;
|
|
}
|
|
|
|
bool HsaBlitManager::hsaCopy(const void *hostSrc, void *hostDst,
|
|
uint32_t size, bool hostToDev) const {
|
|
|
|
// No allocation is necessary for Full Profile
|
|
hsa_status_t status;
|
|
if (roc_device_.agent_profile() == HSA_PROFILE_FULL) {
|
|
status = hsa_memory_copy(hostDst, hostSrc, size);
|
|
if (status != HSA_STATUS_SUCCESS) {
|
|
LogPrintfError("Hsa copy of data failed with code %d", status);
|
|
}
|
|
return (status == HSA_STATUS_SUCCESS);
|
|
}
|
|
|
|
// Allocate requested size of memory
|
|
size_t align = 0x04;
|
|
bool atomics = false;
|
|
void *hsaBuffer = NULL;
|
|
hsaBuffer = roc_device_.hostAlloc(size, align, false);
|
|
if (hsaBuffer == NULL) {
|
|
LogError("Hsa buffer allocation failed with code");
|
|
return false;
|
|
}
|
|
|
|
const hsa_signal_value_t kInitVal = 1;
|
|
hsa_signal_store_relaxed(completion_signal_, kInitVal);
|
|
|
|
// Copy data from Host to Device
|
|
if (hostToDev) {
|
|
memcpy(hsaBuffer, hostSrc, size);
|
|
status = hsa_amd_memory_async_copy(
|
|
hostDst, roc_device_.getBackendDevice(), hsaBuffer,
|
|
roc_device_.getCpuAgent(), size, 0, NULL, completion_signal_);
|
|
if (status == HSA_STATUS_SUCCESS) {
|
|
hsa_signal_value_t val =
|
|
hsa_signal_wait_acquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0,
|
|
uint64_t(-1), HSA_WAIT_STATE_ACTIVE);
|
|
|
|
if (val != (kInitVal - 1)) {
|
|
LogError("Async copy failed");
|
|
status = HSA_STATUS_ERROR;
|
|
}
|
|
}
|
|
else {
|
|
LogPrintfError("Hsa copy from host to device failed with code %d", status);
|
|
}
|
|
|
|
roc_device_.hostFree(hsaBuffer, size);
|
|
return (status == HSA_STATUS_SUCCESS);
|
|
}
|
|
|
|
// Copy data from Device to Host
|
|
status = hsa_amd_memory_async_copy(hsaBuffer, roc_device_.getCpuAgent(),
|
|
hostSrc, roc_device_.getBackendDevice(),
|
|
size, 0, NULL, completion_signal_);
|
|
if (status == HSA_STATUS_SUCCESS) {
|
|
hsa_signal_value_t val = hsa_signal_wait_acquire(
|
|
completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1),
|
|
HSA_WAIT_STATE_ACTIVE);
|
|
|
|
if (val != (kInitVal - 1)) {
|
|
LogError("Async copy failed");
|
|
status = HSA_STATUS_ERROR;
|
|
}
|
|
|
|
if (status == HSA_STATUS_SUCCESS) {
|
|
memcpy(hostDst, hsaBuffer, size);
|
|
}
|
|
} else {
|
|
LogPrintfError("Hsa copy from device to host failed with code %d", status);
|
|
}
|
|
|
|
roc_device_.hostFree(hsaBuffer, size);
|
|
return (status == HSA_STATUS_SUCCESS);
|
|
}
|
|
|
|
bool HsaBlitManager::readBuffer(device::Memory& srcMemory, void* dstHost,
|
|
const amd::Coord3D& origin,
|
|
const amd::Coord3D& size, bool entire) const {
|
|
hsa_memory_register(dstHost, size[0]);
|
|
void* src = static_cast<roc::Memory&>(srcMemory).getDeviceMemory();
|
|
|
|
// Copy data from device to host
|
|
const void *srcDev = reinterpret_cast<const_address>(src) + origin[0];
|
|
bool retval = hsaCopy(srcDev, dstHost, size[0], false);
|
|
|
|
hsa_memory_deregister(dstHost, size[0]);
|
|
return retval;
|
|
}
|
|
|
|
bool HsaBlitManager::readBufferRect(device::Memory& srcMemory, void* dst,
|
|
const amd::BufferRect& bufRect,
|
|
const amd::BufferRect& hostRect,
|
|
const amd::Coord3D& size,
|
|
bool entire) const {
|
|
void* src = static_cast<roc::Memory&>(srcMemory).getDeviceMemory();
|
|
|
|
size_t srcOffset;
|
|
size_t dstOffset;
|
|
|
|
for (size_t z = 0; z < size[2]; ++z) {
|
|
for (size_t y = 0; y < size[1]; ++y) {
|
|
srcOffset = bufRect.offset(0, y, z);
|
|
dstOffset = hostRect.offset(0, y, z);
|
|
|
|
// Copy data from device to host - line by line
|
|
void *dstHost = reinterpret_cast<address>(dst) + dstOffset;
|
|
const void *srcDev = reinterpret_cast<const_address>(src) + srcOffset;
|
|
bool retval = hsaCopy(srcDev, dstHost, size[0], false);
|
|
if (!retval) {
|
|
return retval;
|
|
}
|
|
}
|
|
}
|
|
|
|
return true;
|
|
}
|
|
|
|
static bool hsaCopyImageToBuffer(hsa_agent_t agent,
|
|
hsa_ext_image_t srcImage,
|
|
void* dstBuffer, const amd::Coord3D& srcOrigin,
|
|
const amd::Coord3D& dstOrigin,
|
|
const amd::Coord3D& size, bool entire,
|
|
size_t rowPitch, size_t slicePitch) {
|
|
hsa_ext_image_region_t image_region;
|
|
image_region.offset.x = srcOrigin[0];
|
|
image_region.offset.y = srcOrigin[1];
|
|
image_region.offset.z = srcOrigin[2];
|
|
image_region.range.x = size[0];
|
|
image_region.range.y = size[1];
|
|
image_region.range.z = size[2];
|
|
|
|
char *dstHost = ((char*)dstBuffer) + dstOrigin[0];
|
|
|
|
hsa_status_t status = hsa_ext_image_export(agent, srcImage, dstHost, rowPitch,
|
|
slicePitch, &image_region);
|
|
return (status == HSA_STATUS_SUCCESS);
|
|
}
|
|
|
|
bool HsaBlitManager::readImage(device::Memory& srcMemory, void* dstHost,
|
|
const amd::Coord3D& origin,
|
|
const amd::Coord3D& size, size_t rowPitch,
|
|
size_t slicePitch, bool entire) const {
|
|
roc::Image* srcImage = (roc::Image*)&srcMemory;
|
|
|
|
void* svmDstHost = NULL;
|
|
size_t pinSize = 0;
|
|
FindPinSize(pinSize, size, rowPitch, slicePitch, *srcImage);
|
|
|
|
hsa_agent_t agent = gpu().gpu_device();
|
|
|
|
hsa_status_t status = hsa_amd_memory_lock(dstHost, pinSize,
|
|
&agent, 1, &svmDstHost);
|
|
|
|
if (status != HSA_STATUS_SUCCESS) {
|
|
return false;
|
|
}
|
|
|
|
bool retval = hsaCopyImageToBuffer(agent, srcImage->getHsaImageObject(),
|
|
svmDstHost, origin, amd::Coord3D(0), size, entire,
|
|
rowPitch, slicePitch);
|
|
hsa_amd_memory_unlock(dstHost);
|
|
return retval;
|
|
}
|
|
|
|
bool HsaBlitManager::writeBuffer(const void* srcHost, device::Memory& dstMemory,
|
|
const amd::Coord3D& origin,
|
|
const amd::Coord3D& size, bool entire) const {
|
|
hsa_memory_register(const_cast<void*>(srcHost), size[0]);
|
|
void* dst = static_cast<roc::Memory&>(dstMemory).getDeviceMemory();
|
|
|
|
// Copy data from host to device
|
|
void *dstDev = reinterpret_cast<address>(dst) + origin[0];
|
|
bool retval = hsaCopy(srcHost, dstDev, size[0], true);
|
|
|
|
hsa_memory_deregister(const_cast<void*>(srcHost), size[0]);
|
|
return retval;
|
|
}
|
|
|
|
bool HsaBlitManager::writeBufferRect(const void* src,
|
|
device::Memory& dstMemory,
|
|
const amd::BufferRect& hostRect,
|
|
const amd::BufferRect& bufRect,
|
|
const amd::Coord3D& size,
|
|
bool entire) const {
|
|
void* dst = static_cast<roc::Memory&>(dstMemory).getDeviceMemory();
|
|
|
|
size_t srcOffset;
|
|
size_t dstOffset;
|
|
|
|
for (size_t z = 0; z < size[2]; ++z) {
|
|
for (size_t y = 0; y < size[1]; ++y) {
|
|
srcOffset = hostRect.offset(0, y, z);
|
|
dstOffset = bufRect.offset(0, y, z);
|
|
|
|
// Copy data from host to device - line by line
|
|
void *dstDev = reinterpret_cast<address>(dst) + dstOffset;
|
|
const void *srcHost = reinterpret_cast<const_address>(src) + srcOffset;
|
|
bool retval = hsaCopy(srcHost, dstDev, size[0], true);
|
|
if (!retval) {
|
|
return retval;
|
|
}
|
|
}
|
|
}
|
|
|
|
return true;
|
|
}
|
|
|
|
bool hsaCopyBufferToImage(hsa_agent_t agent, const void* srcBuffer,
|
|
hsa_ext_image_t dstImage,
|
|
const amd::Coord3D& srcOrigin,
|
|
const amd::Coord3D& dstOrigin,
|
|
const amd::Coord3D& size, bool entire,
|
|
size_t rowPitch, size_t slicePitch) {
|
|
char* srcHost = ((char*)srcBuffer) + srcOrigin[0];
|
|
|
|
hsa_ext_image_region_t image_region;
|
|
image_region.offset.x = dstOrigin[0];
|
|
image_region.offset.y = dstOrigin[1];
|
|
image_region.offset.z = dstOrigin[2];
|
|
image_region.range.x = size[0];
|
|
image_region.range.y = size[1];
|
|
image_region.range.z = size[2];
|
|
|
|
hsa_status_t status = hsa_ext_image_import(
|
|
agent, srcHost, rowPitch, slicePitch, dstImage, &image_region);
|
|
return (status == HSA_STATUS_SUCCESS);
|
|
}
|
|
|
|
bool HsaBlitManager::writeImage(const void* srcHost, device::Memory& dstMemory,
|
|
const amd::Coord3D& origin,
|
|
const amd::Coord3D& size, size_t rowPitch,
|
|
size_t slicePitch, bool entire) const {
|
|
roc::Image* image = (roc::Image*)&dstMemory;
|
|
|
|
void* svmSrcHost = NULL;
|
|
size_t pinSize = 0;
|
|
FindPinSize(pinSize, size, rowPitch, slicePitch, *image);
|
|
|
|
hsa_agent_t agent = gpu().gpu_device();
|
|
|
|
hsa_status_t status = hsa_amd_memory_lock(const_cast<void*>(srcHost), pinSize,
|
|
&agent, 1, &svmSrcHost);
|
|
|
|
if (status != HSA_STATUS_SUCCESS) {
|
|
return false;
|
|
}
|
|
|
|
bool retval = hsaCopyBufferToImage(agent, svmSrcHost,
|
|
image->getHsaImageObject(), amd::Coord3D(0),
|
|
origin, size, entire, rowPitch, slicePitch);
|
|
|
|
hsa_amd_memory_unlock(const_cast<void*>(srcHost));
|
|
|
|
return retval;
|
|
}
|
|
|
|
bool HsaBlitManager::copyBuffer(device::Memory& srcMemory,
|
|
device::Memory& dstMemory,
|
|
const amd::Coord3D& srcOrigin,
|
|
const amd::Coord3D& dstOrigin,
|
|
const amd::Coord3D& size, bool entire) const {
|
|
void* src = static_cast<roc::Memory&>(srcMemory).getDeviceMemory();
|
|
void* dst = static_cast<roc::Memory&>(dstMemory).getDeviceMemory();
|
|
|
|
if (srcMemory.isHostMemDirectAccess() && dstMemory.isHostMemDirectAccess()) {
|
|
if (srcMemory.owner()->getMemFlags() & CL_MEM_USE_HOST_PTR) {
|
|
src = srcMemory.owner()->getHostMem();
|
|
}
|
|
|
|
if (dstMemory.owner()->getMemFlags() & CL_MEM_USE_HOST_PTR) {
|
|
dst = dstMemory.owner()->getHostMem();
|
|
}
|
|
}
|
|
|
|
const hsa_agent_t src_agent = (srcMemory.isHostMemDirectAccess())
|
|
? roc_device_.getCpuAgent()
|
|
: roc_device_.getBackendDevice();
|
|
|
|
const hsa_agent_t dst_agent = (dstMemory.isHostMemDirectAccess())
|
|
? roc_device_.getCpuAgent()
|
|
: roc_device_.getBackendDevice();
|
|
|
|
// Straight forward buffer copy
|
|
const hsa_signal_value_t kInitVal = 1;
|
|
hsa_signal_store_relaxed(completion_signal_, kInitVal);
|
|
hsa_status_t status = hsa_amd_memory_async_copy(
|
|
(reinterpret_cast<address>(dst) + dstOrigin[0]), dst_agent,
|
|
(reinterpret_cast<const_address>(src) + srcOrigin[0]), src_agent, size[0],
|
|
0, NULL, completion_signal_);
|
|
if (status != HSA_STATUS_SUCCESS) {
|
|
LogPrintfError("DMA buffer failed with code %d", status);
|
|
return false;
|
|
}
|
|
|
|
hsa_signal_value_t val =
|
|
hsa_signal_wait_acquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0,
|
|
uint64_t(-1), HSA_WAIT_STATE_ACTIVE);
|
|
|
|
if (val != (kInitVal - 1)) {
|
|
LogError("Async copy failed");
|
|
return false;
|
|
}
|
|
|
|
return true;
|
|
}
|
|
|
|
bool HsaBlitManager::copyBufferRect(device::Memory& srcMemory,
|
|
device::Memory& dstMemory,
|
|
const amd::BufferRect& srcRect,
|
|
const amd::BufferRect& dstRect,
|
|
const amd::Coord3D& size,
|
|
bool entire) const {
|
|
void* src = static_cast<roc::Memory&>(srcMemory).getDeviceMemory();
|
|
void* dst = static_cast<roc::Memory&>(dstMemory).getDeviceMemory();
|
|
|
|
const hsa_signal_value_t kInitVal = size[2] * size[1];
|
|
hsa_signal_store_relaxed(completion_signal_, kInitVal);
|
|
|
|
for (size_t z = 0; z < size[2]; ++z) {
|
|
for (size_t y = 0; y < size[1]; ++y) {
|
|
size_t srcOffset = srcRect.offset(0, y, z);
|
|
size_t dstOffset = dstRect.offset(0, y, z);
|
|
|
|
// Copy memory line by line
|
|
hsa_status_t status = hsa_amd_memory_async_copy(
|
|
(reinterpret_cast<address>(dst) + dstOffset),
|
|
roc_device_.getBackendDevice(),
|
|
(reinterpret_cast<const_address>(src) + srcOffset),
|
|
roc_device_.getBackendDevice(), size[0], 0, NULL,
|
|
completion_signal_);
|
|
if (status != HSA_STATUS_SUCCESS) {
|
|
LogPrintfError("DMA buffer failed with code %d", status);
|
|
return false;
|
|
}
|
|
}
|
|
}
|
|
|
|
hsa_signal_value_t val =
|
|
hsa_signal_wait_acquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ,
|
|
0, uint64_t(-1), HSA_WAIT_STATE_ACTIVE);
|
|
|
|
if (val != 0) {
|
|
LogError("Async copy failed");
|
|
return false;
|
|
}
|
|
|
|
return true;
|
|
}
|
|
|
|
bool HsaBlitManager::copyImageToBuffer(device::Memory& srcMemory,
|
|
device::Memory& dstMemory,
|
|
const amd::Coord3D& srcOrigin,
|
|
const amd::Coord3D& dstOrigin,
|
|
const amd::Coord3D& size, bool entire,
|
|
size_t rowPitch,
|
|
size_t slicePitch) const {
|
|
roc::Image& srcImage = (roc::Image&)srcMemory;
|
|
roc::Buffer& dstBuffer = (roc::Buffer&)dstMemory;
|
|
|
|
return hsaCopyImageToBuffer(gpu().gpu_device(), srcImage.getHsaImageObject(),
|
|
dstBuffer.getDeviceMemory(), srcOrigin, dstOrigin,
|
|
size, entire, rowPitch, slicePitch);
|
|
}
|
|
|
|
bool HsaBlitManager::copyBufferToImage(device::Memory& srcMemory,
|
|
device::Memory& dstMemory,
|
|
const amd::Coord3D& srcOrigin,
|
|
const amd::Coord3D& dstOrigin,
|
|
const amd::Coord3D& size, bool entire,
|
|
size_t rowPitch,
|
|
size_t slicePitch) const {
|
|
roc::Buffer& srcBuffer = (roc::Buffer&)srcMemory;
|
|
roc::Image& dstImage = (roc::Image&)dstMemory;
|
|
|
|
return hsaCopyBufferToImage(gpu().gpu_device(), srcBuffer.getDeviceMemory(),
|
|
dstImage.getHsaImageObject(), srcOrigin,
|
|
dstOrigin, size, entire, rowPitch, slicePitch);
|
|
}
|
|
|
|
bool HsaBlitManager::copyImage(device::Memory& srcMemory,
|
|
device::Memory& dstMemory,
|
|
const amd::Coord3D& srcOrigin,
|
|
const amd::Coord3D& dstOrigin,
|
|
const amd::Coord3D& size, bool entire) const {
|
|
if (srcMemory.isHostMemDirectAccess() &&
|
|
dstMemory.isHostMemDirectAccess()) {
|
|
return device::HostBlitManager::copyImage(
|
|
srcMemory, dstMemory, srcOrigin, dstOrigin, size, entire);
|
|
}
|
|
|
|
roc::Image *srcImage = (roc::Image *)&srcMemory;
|
|
roc::Image *dstImage = (roc::Image *)&dstMemory;
|
|
|
|
hsa_dim3_t src_offset = { 0 };
|
|
src_offset.x = srcOrigin[0];
|
|
src_offset.y = srcOrigin[1];
|
|
src_offset.z = srcOrigin[2];
|
|
|
|
hsa_dim3_t dst_offset = { 0 };
|
|
dst_offset.x = dstOrigin[0];
|
|
dst_offset.y = dstOrigin[1];
|
|
dst_offset.z = dstOrigin[2];
|
|
|
|
hsa_dim3_t copy_size = { 0 };
|
|
copy_size.x = size[0];
|
|
copy_size.y = size[1];
|
|
copy_size.z = size[2];
|
|
|
|
hsa_status_t status = hsa_ext_image_copy(
|
|
gpu().gpu_device(), srcImage->getHsaImageObject(), &src_offset,
|
|
dstImage->getHsaImageObject(), &dst_offset, ©_size);
|
|
return (status == HSA_STATUS_SUCCESS);
|
|
}
|
|
|
|
bool HsaBlitManager::fillBuffer(device::Memory& memory, const void* pattern,
|
|
size_t patternSize, const amd::Coord3D& origin,
|
|
const amd::Coord3D& size, bool entire) const {
|
|
void* fillMem = static_cast<roc::Memory&>(memory).getDeviceMemory();
|
|
|
|
size_t offset = origin[0];
|
|
size_t fillSize = size[0];
|
|
|
|
if ((fillSize % patternSize) != 0) {
|
|
LogError("Misaligned buffer size and pattern size!");
|
|
}
|
|
|
|
// Fill the buffer memory with a pattern
|
|
for (size_t i = 0; i < (fillSize / patternSize); i++) {
|
|
void *dstDev = reinterpret_cast<address>(fillMem) + offset;
|
|
bool retval = hsaCopy(pattern, dstDev, patternSize, true);
|
|
if (!retval) {
|
|
LogError("DMA buffer failed with code");
|
|
return retval;
|
|
}
|
|
|
|
offset += patternSize;
|
|
}
|
|
|
|
return true;
|
|
}
|
|
|
|
bool HsaBlitManager::fillImage(device::Memory& memory, const void* pattern,
|
|
const amd::Coord3D& origin,
|
|
const amd::Coord3D& size, bool entire) const {
|
|
if (memory.isHostMemDirectAccess()) {
|
|
return device::HostBlitManager::fillImage(memory, pattern, origin, size, entire);
|
|
}
|
|
|
|
roc::Image *image = (roc::Image*)&memory;
|
|
hsa_ext_image_region_t image_region;
|
|
image_region.offset.x = origin[0];
|
|
image_region.offset.y = origin[1];
|
|
image_region.offset.z = origin[2];
|
|
image_region.range.x = size[0];
|
|
image_region.range.y = size[1];
|
|
image_region.range.z = size[2];
|
|
|
|
hsa_status_t status = hsa_ext_image_clear(
|
|
gpu().gpu_device(), image->getHsaImageObject(),
|
|
pattern, &image_region);
|
|
return (status == HSA_STATUS_SUCCESS);
|
|
}
|
|
|
|
static void
|
|
CalcRowSlicePitches(
|
|
cl_ulong* pitch, const cl_int* copySize,
|
|
size_t rowPitch, size_t slicePitch, const Memory& mem)
|
|
{
|
|
const roc::Image &hsaImage = static_cast< const roc::Image &>(mem);
|
|
bool img1Darray =
|
|
(mem.owner()->getType() == CL_MEM_OBJECT_IMAGE1D_ARRAY) ? true : false;
|
|
size_t memFmtSize = mem.owner()->asImage()->getImageFormat().getElementSize();
|
|
|
|
if (rowPitch == 0) {
|
|
pitch[0] = copySize[0];
|
|
}
|
|
else {
|
|
pitch[0] = rowPitch / memFmtSize;
|
|
}
|
|
if (slicePitch == 0) {
|
|
pitch[1] = pitch[0] * (img1Darray ? 1 : copySize[1]);
|
|
}
|
|
else {
|
|
pitch[1] = slicePitch / memFmtSize;
|
|
}
|
|
assert((pitch[0] <= pitch[1]) && "rowPitch must be <= slicePitch");
|
|
|
|
if (img1Darray) {
|
|
// For 1D array rowRitch = slicePitch
|
|
pitch[0] = pitch[1];
|
|
}
|
|
}
|
|
|
|
KernelBlitManager::KernelBlitManager(device::VirtualDevice& vDev, Setup setup)
|
|
: HsaBlitManager(vDev, setup),
|
|
context_(NULL),
|
|
program_(NULL)
|
|
{
|
|
for (uint i = 0; i < BlitTotal; ++i) {
|
|
kernels_[i] = NULL;
|
|
}
|
|
}
|
|
|
|
KernelBlitManager::~KernelBlitManager()
|
|
{
|
|
for (uint i = 0; i < BlitTotal; ++i) {
|
|
if (NULL != kernels_[i]) {
|
|
kernels_[i]->release();
|
|
}
|
|
}
|
|
|
|
if (NULL != program_) {
|
|
program_->release();
|
|
}
|
|
|
|
if (NULL != context_) {
|
|
// Release a dummy context
|
|
context_->release();
|
|
}
|
|
}
|
|
|
|
bool
|
|
KernelBlitManager::readBuffer(
|
|
device::Memory& srcMemory,
|
|
void* dstHost,
|
|
const amd::Coord3D& origin,
|
|
const amd::Coord3D& size,
|
|
bool entire) const
|
|
{
|
|
//if (setup_.disableReadBuffer_ || srcMemory.isHostMemDirectAccess()) {
|
|
// return device::HostBlitManager::readBuffer(srcMemory, dstHost, origin,
|
|
// size, entire);
|
|
//}
|
|
// Exercise HSA path for now.
|
|
return HsaBlitManager::readBuffer(srcMemory, dstHost, origin,
|
|
size, entire);
|
|
|
|
amd::Buffer *dstMemory = new (*context_) amd::Buffer(
|
|
*context_, CL_MEM_USE_HOST_PTR, size[0]);
|
|
|
|
if (!dstMemory->create(const_cast<void *>(dstHost))) {
|
|
LogError("[OCL] Fail to create mem object for destination");
|
|
return false;
|
|
}
|
|
|
|
device::Memory *devDstMemory = dstMemory->getDeviceMemory(dev_);
|
|
if (devDstMemory== NULL) {
|
|
LogError("[OCL] Fail to create device mem object for destination");
|
|
return false;
|
|
}
|
|
|
|
bool result = copyBuffer(
|
|
srcMemory, *devDstMemory, origin, amd::Coord3D(0), size, entire);
|
|
|
|
// Wait for the transfer to finish so that we could safely release the
|
|
// destination memory object.
|
|
// TODO: we could remove this if issue on implicit memory registration is
|
|
// fixed by KFD, so that we could pass the pattern as SVM.
|
|
gpu().releaseGpuMemoryFence();
|
|
|
|
dstMemory->release();
|
|
|
|
return result;
|
|
}
|
|
|
|
bool
|
|
KernelBlitManager::readBufferRect(
|
|
device::Memory& srcMemory,
|
|
void* dstHost,
|
|
const amd::BufferRect& bufRect,
|
|
const amd::BufferRect& hostRect,
|
|
const amd::Coord3D& size,
|
|
bool entire) const
|
|
{
|
|
// if (setup_.disableReadBufferRect_ || srcMemory.isHostMemDirectAccess()) {
|
|
//return device::HostBlitManager::readBufferRect(
|
|
// srcMemory, dstHost, bufRect, hostRect, size, entire);
|
|
// }
|
|
|
|
// Exercise HSA path for now.
|
|
return HsaBlitManager::readBufferRect(
|
|
srcMemory, dstHost, bufRect, hostRect, size, entire);
|
|
|
|
size_t dstSize = hostRect.start_ + hostRect.end_;
|
|
amd::Buffer *dstMemory =
|
|
new (*context_) amd::Buffer(*context_, CL_MEM_USE_HOST_PTR, dstSize);
|
|
|
|
if (!dstMemory->create(const_cast<void *>(dstHost))) {
|
|
LogError("[OCL] Fail to create mem object for destination");
|
|
return false;
|
|
}
|
|
|
|
device::Memory *devDstMemory = dstMemory->getDeviceMemory(dev_);
|
|
if (devDstMemory== NULL) {
|
|
LogError("[OCL] Fail to create device mem object for destination");
|
|
return false;
|
|
}
|
|
|
|
bool result = copyBufferRect(
|
|
srcMemory, *devDstMemory, bufRect, hostRect, size, entire);
|
|
|
|
// Wait for the transfer to finish so that we could safely release the
|
|
// destination memory object.
|
|
// TODO: we could remove this if issue on implicit memory registration is
|
|
// fixed by KFD, so that we could pass the pattern as SVM.
|
|
gpu().releaseGpuMemoryFence();
|
|
|
|
dstMemory->release();
|
|
|
|
return result;
|
|
}
|
|
|
|
void
|
|
FindLinearSize(
|
|
size_t& linearSize, const amd::Coord3D& size,
|
|
size_t& rowPitch, size_t& slicePitch, const device::Memory& mem)
|
|
{
|
|
const roc::Image &image = static_cast<const roc::Image &>(mem);
|
|
size_t elementSize = mem.owner()->asImage()->getImageFormat().getElementSize();
|
|
|
|
linearSize = size[0] * elementSize;
|
|
if ((rowPitch == 0) || (rowPitch == linearSize)) {
|
|
rowPitch = 0;
|
|
}
|
|
else {
|
|
linearSize = rowPitch;
|
|
}
|
|
|
|
// Calculate the pin size, which should be equal to the copy size
|
|
for (uint i = 1; i < mem.owner()->asImage()->getDims(); ++i) {
|
|
linearSize *= size[i];
|
|
if (i == 1) {
|
|
if ((slicePitch == 0) || (slicePitch == linearSize)) {
|
|
slicePitch = 0;
|
|
}
|
|
else {
|
|
if (mem.owner()->getType() != CL_MEM_OBJECT_IMAGE1D_ARRAY) {
|
|
linearSize = slicePitch;
|
|
}
|
|
else {
|
|
linearSize = slicePitch * size[i];
|
|
}
|
|
}
|
|
}
|
|
}
|
|
}
|
|
|
|
// The following data structures will be used for the view creations.
|
|
// Some formats has to be converted before a kernel blit operation
|
|
struct FormatConvertion {
|
|
cl_uint clOldType_;
|
|
cl_uint clNewType_;
|
|
};
|
|
|
|
// The list of rejected data formats and corresponding conversion
|
|
static const FormatConvertion RejectedData[] =
|
|
{
|
|
{ CL_UNORM_INT8, CL_UNSIGNED_INT8 },
|
|
{ CL_UNORM_INT16, CL_UNSIGNED_INT16 },
|
|
{ CL_SNORM_INT8, CL_UNSIGNED_INT8 },
|
|
{ CL_SNORM_INT16, CL_UNSIGNED_INT16 },
|
|
{ CL_HALF_FLOAT, CL_UNSIGNED_INT16 },
|
|
{ CL_FLOAT, CL_UNSIGNED_INT32 },
|
|
{ CL_SIGNED_INT8, CL_UNSIGNED_INT8 },
|
|
{ CL_SIGNED_INT16, CL_UNSIGNED_INT16 },
|
|
{ CL_UNORM_INT_101010, CL_UNSIGNED_INT8 },
|
|
{ CL_SIGNED_INT32, CL_UNSIGNED_INT32 }
|
|
};
|
|
|
|
// The list of rejected channel's order and corresponding conversion
|
|
static const FormatConvertion RejectedOrder[] =
|
|
{
|
|
{ CL_A, CL_R },
|
|
{ CL_RA, CL_RG },
|
|
{ CL_LUMINANCE, CL_R },
|
|
{ CL_INTENSITY, CL_R },
|
|
{ CL_RGB, CL_RGBA },
|
|
{ CL_BGRA, CL_RGBA },
|
|
{ CL_ARGB, CL_RGBA },
|
|
{ CL_sRGB, CL_RGBA },
|
|
{ CL_sRGBx, CL_RGBA },
|
|
{ CL_sRGBA, CL_RGBA },
|
|
{ CL_sBGRA, CL_RGBA },
|
|
{ CL_DEPTH, CL_R}
|
|
};
|
|
|
|
const uint RejectedFormatDataTotal =
|
|
sizeof(RejectedData) / sizeof(FormatConvertion);
|
|
const uint RejectedFormatChannelTotal =
|
|
sizeof(RejectedOrder) / sizeof(FormatConvertion);
|
|
|
|
amd::Image::Format
|
|
KernelBlitManager::filterFormat(amd::Image::Format oldFormat) const
|
|
{
|
|
cl_image_format newFormat;
|
|
newFormat.image_channel_data_type = oldFormat.image_channel_data_type;
|
|
newFormat.image_channel_order = oldFormat.image_channel_order;
|
|
|
|
// Find unsupported formats
|
|
for (uint i = 0; i < RejectedFormatDataTotal; ++i) {
|
|
if (RejectedData[i].clOldType_ == oldFormat.image_channel_data_type) {
|
|
newFormat.image_channel_data_type = RejectedData[i].clNewType_;
|
|
break;
|
|
}
|
|
}
|
|
|
|
// Find unsupported channel's order
|
|
for (uint i = 0; i < RejectedFormatChannelTotal; ++i) {
|
|
if (RejectedOrder[i].clOldType_ == oldFormat.image_channel_order) {
|
|
newFormat.image_channel_order = RejectedOrder[i].clNewType_;
|
|
break;
|
|
}
|
|
}
|
|
|
|
return amd::Image::Format(newFormat);
|
|
}
|
|
|
|
device::Memory *
|
|
KernelBlitManager::createImageView(
|
|
device::Memory &parent,
|
|
amd::Image::Format newFormat) const
|
|
{
|
|
amd::Image *image =
|
|
parent.owner()->asImage()->createView(parent.owner()->getContext(), newFormat, &gpu());
|
|
|
|
if (image == NULL) {
|
|
LogError("[OCL] Fail to allocate view of image object");
|
|
return NULL;
|
|
}
|
|
|
|
Image* devImage = new roc::Image(static_cast<const Device &>(dev_), *image);
|
|
if (devImage == NULL) {
|
|
LogError("[OCL] Fail to allocate device mem object for the view");
|
|
image->release();
|
|
return NULL;
|
|
}
|
|
|
|
if (!devImage->createView(static_cast<roc::Image &>(parent))) {
|
|
LogError("[OCL] Fail to create device mem object for the view");
|
|
delete devImage;
|
|
image->release();
|
|
return NULL;
|
|
}
|
|
|
|
image->replaceDeviceMemory(&dev_, devImage);
|
|
|
|
return devImage;
|
|
}
|
|
|
|
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
|
|
{
|
|
return HsaBlitManager::readImage(
|
|
srcMemory, dstHost, origin, size, rowPitch, slicePitch, entire);
|
|
}
|
|
|
|
bool
|
|
KernelBlitManager::writeBuffer(
|
|
const void* srcHost,
|
|
device::Memory& dstMemory,
|
|
const amd::Coord3D& origin,
|
|
const amd::Coord3D& size,
|
|
bool entire) const
|
|
{
|
|
// if (setup_.disableWriteBuffer_ || dstMemory.isHostMemDirectAccess()) {
|
|
//return device::HostBlitManager::writeBuffer(srcHost, dstMemory, origin, size,
|
|
// entire);
|
|
// }
|
|
|
|
// Exercise HSA path for now.
|
|
return HsaBlitManager::writeBuffer(srcHost, dstMemory, origin, size,
|
|
entire);
|
|
|
|
amd::Buffer *srcMemory =
|
|
new (*context_) amd::Buffer(*context_, CL_MEM_USE_HOST_PTR, size[0]);
|
|
|
|
if (!srcMemory->create(const_cast<void *>(srcHost))) {
|
|
LogError("[OCL] Fail to create mem object for destination");
|
|
return false;
|
|
}
|
|
|
|
device::Memory *devSrcMemory = srcMemory->getDeviceMemory(dev_);
|
|
if (devSrcMemory== NULL) {
|
|
LogError("[OCL] Fail to create device mem object for destination");
|
|
return false;
|
|
}
|
|
|
|
bool result =
|
|
copyBuffer(*devSrcMemory, dstMemory, amd::Coord3D(0), origin, size, entire);
|
|
|
|
// Wait for the transfer to finish so that we could safely release the
|
|
// source memory object.
|
|
// TODO: we could remove this if issue on implicit memory registration is
|
|
// fixed by KFD, so that we could pass the pattern as SVM.
|
|
gpu().releaseGpuMemoryFence();
|
|
|
|
srcMemory->release();
|
|
|
|
return result;
|
|
}
|
|
|
|
bool
|
|
KernelBlitManager::writeBufferRect(
|
|
const void* srcHost,
|
|
device::Memory& dstMemory,
|
|
const amd::BufferRect& hostRect,
|
|
const amd::BufferRect& bufRect,
|
|
const amd::Coord3D& size,
|
|
bool entire) const
|
|
{
|
|
// if (setup_.disableWriteBufferRect_ || dstMemory.isHostMemDirectAccess()) {
|
|
//return device::HostBlitManager::writeBufferRect(
|
|
// srcHost, dstMemory, hostRect, bufRect, size, entire);
|
|
// }
|
|
|
|
// Exercise HSA path for now.
|
|
return HsaBlitManager::writeBufferRect(
|
|
srcHost, dstMemory, hostRect, bufRect, size, entire);
|
|
|
|
size_t srcSize = hostRect.start_ + hostRect.end_;
|
|
amd::Buffer *srcMemory =
|
|
new (*context_) amd::Buffer(*context_, CL_MEM_USE_HOST_PTR, srcSize);
|
|
|
|
if (!srcMemory->create(const_cast<void *>(srcHost))) {
|
|
LogError("[OCL] Fail to create mem object for destination");
|
|
return false;
|
|
}
|
|
|
|
device::Memory *devSrcMemory = srcMemory->getDeviceMemory(dev_);
|
|
if (devSrcMemory== NULL) {
|
|
LogError("[OCL] Fail to create device mem object for destination");
|
|
return false;
|
|
}
|
|
|
|
bool result = copyBufferRect(
|
|
*devSrcMemory, dstMemory, hostRect, bufRect, size, entire);
|
|
|
|
// Wait for the transfer to finish so that we could safely release the
|
|
// destination memory object.
|
|
// TODO: we could remove this if issue on implicit memory registration is
|
|
// fixed by KFD, so that we could pass the pattern as SVM.
|
|
gpu().releaseGpuMemoryFence();
|
|
|
|
srcMemory->release();
|
|
|
|
return result;
|
|
}
|
|
|
|
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
|
|
{
|
|
return HsaBlitManager::writeImage(
|
|
srcHost, dstMemory, origin, size, rowPitch, slicePitch, entire);
|
|
}
|
|
|
|
bool
|
|
KernelBlitManager::copyBuffer(
|
|
device::Memory& srcMemory,
|
|
device::Memory& dstMemory,
|
|
const amd::Coord3D& srcOrigin,
|
|
const amd::Coord3D& dstOrigin,
|
|
const amd::Coord3D& sizeIn,
|
|
bool entire) const
|
|
{
|
|
// if (setup_.disableCopyBuffer_ ||
|
|
// (srcMemory.isHostMemDirectAccess() &&
|
|
// dstMemory.isHostMemDirectAccess())) {
|
|
//return HsaBlitManager::copyBuffer(
|
|
// srcMemory, dstMemory, srcOrigin, dstOrigin, sizeIn, entire);
|
|
// }
|
|
|
|
// Exercise HSA path for now.
|
|
return HsaBlitManager::copyBuffer(
|
|
srcMemory, dstMemory, srcOrigin, dstOrigin, sizeIn, entire);
|
|
|
|
uint blitType = BlitCopyBuffer;
|
|
size_t dim = 1;
|
|
size_t globalWorkOffset[3] = { 0, 0, 0 };
|
|
size_t globalWorkSize = 0;
|
|
size_t localWorkSize = 0;
|
|
|
|
const static uint CopyBuffAlignment[3] = { 16, 4, 1 };
|
|
amd::Coord3D size(sizeIn[0], sizeIn[1], sizeIn[2]);
|
|
|
|
bool aligned;
|
|
uint i;
|
|
for (i = 0; i < 3; ++i) {
|
|
// Check source alignments
|
|
aligned = ((srcOrigin[0] % CopyBuffAlignment[i]) == 0);
|
|
// Check destination alignments
|
|
aligned &= ((dstOrigin[0] % CopyBuffAlignment[i]) == 0);
|
|
// Check copy size alignment in the first dimension
|
|
aligned &= ((sizeIn[0] % CopyBuffAlignment[i]) == 0);
|
|
|
|
if (aligned) {
|
|
if (CopyBuffAlignment[i] != 1) {
|
|
blitType = BlitCopyBufferAligned;
|
|
}
|
|
break;
|
|
}
|
|
}
|
|
|
|
cl_uint remain;
|
|
if (blitType == BlitCopyBufferAligned) {
|
|
size.c[0] /= CopyBuffAlignment[i];
|
|
}
|
|
else {
|
|
remain = size[0] % 4;
|
|
size.c[0] /= 4;
|
|
size.c[0] += 1;
|
|
}
|
|
|
|
// Program the dispatch dimensions
|
|
localWorkSize = 256;
|
|
globalWorkSize = amd::alignUp(size[0] , 256);
|
|
|
|
// Program kernels arguments for the blit operation
|
|
cl_mem clmem = ((cl_mem) as_cl<amd::Memory>(srcMemory.owner()));
|
|
kernels_[blitType]->parameters().set(0, sizeof(cl_mem), &clmem);
|
|
clmem = ((cl_mem) as_cl<amd::Memory>(dstMemory.owner()));
|
|
kernels_[blitType]->parameters().set(1, sizeof(cl_mem), &clmem);
|
|
// Program source origin
|
|
cl_ulong srcOffset = srcOrigin[0] / CopyBuffAlignment[i];
|
|
kernels_[blitType]->parameters().set(2, sizeof(srcOffset), &srcOffset);
|
|
|
|
// Program destination origin
|
|
cl_ulong dstOffset = dstOrigin[0] / CopyBuffAlignment[i];
|
|
kernels_[blitType]->parameters().set(3, sizeof(dstOffset), &dstOffset);
|
|
|
|
cl_ulong copySize = size[0];
|
|
kernels_[blitType]->parameters().set(4, sizeof(copySize), ©Size);
|
|
|
|
if (blitType == BlitCopyBufferAligned) {
|
|
cl_int alignment = CopyBuffAlignment[i];
|
|
kernels_[blitType]->parameters().set(5, sizeof(alignment), &alignment);
|
|
}
|
|
else {
|
|
kernels_[blitType]->parameters().set(5, sizeof(remain), &remain);
|
|
}
|
|
|
|
// Create ND range object for the kernel's execution
|
|
amd::NDRangeContainer ndrange(
|
|
1, globalWorkOffset, &globalWorkSize, &localWorkSize);
|
|
|
|
// Execute the blit
|
|
address parameters = kernels_[blitType]->parameters().capture(dev_);
|
|
bool result = gpu().submitKernelInternal(
|
|
ndrange, *kernels_[blitType], parameters, NULL);
|
|
kernels_[blitType]->parameters().release(const_cast<address>(parameters), dev_);
|
|
return result;
|
|
}
|
|
|
|
bool
|
|
KernelBlitManager::copyBufferRect(
|
|
device::Memory& srcMemory,
|
|
device::Memory& dstMemory,
|
|
const amd::BufferRect& srcRectIn,
|
|
const amd::BufferRect& dstRectIn,
|
|
const amd::Coord3D& sizeIn,
|
|
bool entire) const
|
|
{
|
|
// if (setup_.disableCopyBuffer_ ||
|
|
// (srcMemory.isHostMemDirectAccess() && dstMemory.isHostMemDirectAccess())) {
|
|
//return HsaBlitManager::copyBufferRect(
|
|
// srcMemory, dstMemory, srcRectIn, dstRectIn, sizeIn, entire);
|
|
// }
|
|
|
|
// Exercise HSA path for now.
|
|
return HsaBlitManager::copyBufferRect(
|
|
srcMemory, dstMemory, srcRectIn, dstRectIn, sizeIn, entire);
|
|
|
|
uint blitType = BlitCopyBufferRect;
|
|
size_t dim = 3;
|
|
size_t globalWorkOffset[3] = { 0, 0, 0 };
|
|
size_t globalWorkSize[3];
|
|
size_t localWorkSize[3];
|
|
|
|
const static uint CopyRectAlignment[3] = { 16, 4, 1 };
|
|
|
|
bool aligned;
|
|
uint i;
|
|
for (i = 0; i < sizeof(CopyRectAlignment) / sizeof(uint); i++) {
|
|
// Check source alignments
|
|
aligned = ((srcRectIn.rowPitch_ % CopyRectAlignment[i]) == 0);
|
|
aligned &= ((srcRectIn.slicePitch_ % CopyRectAlignment[i]) == 0);
|
|
aligned &= ((srcRectIn.start_ % CopyRectAlignment[i]) == 0);
|
|
|
|
// Check destination alignments
|
|
aligned &= ((dstRectIn.rowPitch_ % CopyRectAlignment[i]) == 0);
|
|
aligned &= ((dstRectIn.slicePitch_ % CopyRectAlignment[i]) == 0);
|
|
aligned &= ((dstRectIn.start_ % CopyRectAlignment[i]) == 0);
|
|
|
|
// Check copy size alignment in the first dimension
|
|
aligned &= ((sizeIn[0] % CopyRectAlignment[i]) == 0);
|
|
|
|
if (aligned) {
|
|
if (CopyRectAlignment[i] != 1) {
|
|
blitType = BlitCopyBufferRectAligned;
|
|
}
|
|
break;
|
|
}
|
|
}
|
|
|
|
amd::BufferRect srcRect;
|
|
amd::BufferRect dstRect;
|
|
amd::Coord3D size(sizeIn[0], sizeIn[1], sizeIn[2]);
|
|
|
|
srcRect.rowPitch_ = srcRectIn.rowPitch_ / CopyRectAlignment[i];
|
|
srcRect.slicePitch_ = srcRectIn.slicePitch_ / CopyRectAlignment[i];
|
|
srcRect.start_ = srcRectIn.start_ / CopyRectAlignment[i];
|
|
srcRect.end_ = srcRectIn.end_ / CopyRectAlignment[i];
|
|
|
|
dstRect.rowPitch_ = dstRectIn.rowPitch_ / CopyRectAlignment[i];
|
|
dstRect.slicePitch_ = dstRectIn.slicePitch_ / CopyRectAlignment[i];
|
|
dstRect.start_ = dstRectIn.start_ / CopyRectAlignment[i];
|
|
dstRect.end_ = dstRectIn.end_ / CopyRectAlignment[i];
|
|
|
|
size.c[0] /= CopyRectAlignment[i];
|
|
|
|
// Program the kernel's workload depending on the transfer dimensions
|
|
if ((size[1] == 1) && (size[2] == 1)) {
|
|
globalWorkSize[0] = amd::alignUp(size[0], 256);
|
|
globalWorkSize[1] = 1;
|
|
globalWorkSize[2] = 1;
|
|
localWorkSize[0] = 256;
|
|
localWorkSize[1] = 1;
|
|
localWorkSize[2] = 1;
|
|
}
|
|
else if (size[2] == 1) {
|
|
globalWorkSize[0] = amd::alignUp(size[0], 16);
|
|
globalWorkSize[1] = amd::alignUp(size[1], 16);
|
|
globalWorkSize[2] = 1;
|
|
localWorkSize[0] = localWorkSize[1] = 16;
|
|
localWorkSize[2] = 1;
|
|
}
|
|
else {
|
|
globalWorkSize[0] = amd::alignUp(size[0], 8);
|
|
globalWorkSize[1] = amd::alignUp(size[1], 8);
|
|
globalWorkSize[2] = amd::alignUp(size[2], 4);
|
|
localWorkSize[0] = localWorkSize[1] = 8;
|
|
localWorkSize[2] = 4;
|
|
}
|
|
|
|
|
|
// Program kernels arguments for the blit operation
|
|
cl_mem clmem = ((cl_mem) as_cl<amd::Memory>(srcMemory.owner()));
|
|
kernels_[blitType]->parameters().set(0, sizeof(cl_mem), &clmem);
|
|
clmem = ((cl_mem) as_cl<amd::Memory>(dstMemory.owner()));
|
|
kernels_[blitType]->parameters().set(1, sizeof(cl_mem), &clmem);
|
|
cl_ulong src[4] = {srcRect.rowPitch_,
|
|
srcRect.slicePitch_,
|
|
srcRect.start_, 0 };
|
|
kernels_[blitType]->parameters().set(2, sizeof(src), src);
|
|
cl_ulong dst[4] = {dstRect.rowPitch_,
|
|
dstRect.slicePitch_,
|
|
dstRect.start_, 0 };
|
|
kernels_[blitType]->parameters().set(3, sizeof(dst), dst);
|
|
cl_ulong copySize[4] = {size[0],
|
|
size[1],
|
|
size[2],
|
|
CopyRectAlignment[i] };
|
|
kernels_[blitType]->parameters().set(4, sizeof(copySize), copySize);
|
|
|
|
// Create ND range object for the kernel's execution
|
|
amd::NDRangeContainer ndrange(dim,
|
|
globalWorkOffset, globalWorkSize, localWorkSize);
|
|
|
|
// Execute the blit
|
|
address parameters = kernels_[blitType]->parameters().capture(dev_);
|
|
bool result = gpu().submitKernelInternal(
|
|
ndrange, *kernels_[blitType], parameters, NULL);
|
|
kernels_[blitType]->parameters().release(const_cast<address>(parameters), dev_);
|
|
return result;
|
|
}
|
|
|
|
bool
|
|
KernelBlitManager::copyImageToBuffer(
|
|
device::Memory& srcMemory,
|
|
device::Memory& dstMemory,
|
|
const amd::Coord3D& srcOrigin,
|
|
const amd::Coord3D& dstOrigin,
|
|
const amd::Coord3D& size,
|
|
bool entire,
|
|
size_t rowPitch,
|
|
size_t slicePitch) const
|
|
{
|
|
if (dstMemory.isHostMemDirectAccess()) {
|
|
return HsaBlitManager::copyImageToBuffer(srcMemory, dstMemory, srcOrigin,
|
|
dstOrigin, size, entire, rowPitch,
|
|
slicePitch);
|
|
}
|
|
|
|
amd::Image::Format oldFormat = srcMemory.owner()->asImage()->getImageFormat();
|
|
amd::Image::Format newFormat = filterFormat(oldFormat);
|
|
bool useView = false;
|
|
|
|
device::Memory* srcView = &srcMemory;
|
|
if (oldFormat != newFormat) {
|
|
srcView = createImageView(srcMemory, newFormat);
|
|
useView = true;
|
|
}
|
|
|
|
roc::Image& srcImage = static_cast<roc::Image&>(*srcView);
|
|
|
|
amd::Image* image = srcImage.owner()->asImage();
|
|
uint blitType = 0;
|
|
blitType = BlitCopyImageToBuffer;
|
|
size_t dim = 0;
|
|
size_t globalWorkOffset[3] = {0, 0, 0};
|
|
size_t globalWorkSize[3];
|
|
size_t localWorkSize[3];
|
|
|
|
// Program the kernels workload depending on the blit dimensions
|
|
const size_t imageDims = srcImage.owner()->asImage()->getDims();
|
|
dim = 3;
|
|
// Find the current blit type
|
|
if (imageDims == 1) {
|
|
globalWorkSize[0] = amd::alignUp(size[0], 256);
|
|
globalWorkSize[1] = amd::alignUp(size[1], 1);
|
|
globalWorkSize[2] = amd::alignUp(size[2], 1);
|
|
localWorkSize[0] = 256;
|
|
localWorkSize[1] = localWorkSize[2] = 1;
|
|
} else if (imageDims == 2) {
|
|
globalWorkSize[0] = amd::alignUp(size[0], 16);
|
|
globalWorkSize[1] = amd::alignUp(size[1], 16);
|
|
globalWorkSize[2] = amd::alignUp(size[2], 1);
|
|
localWorkSize[0] = localWorkSize[1] = 16;
|
|
localWorkSize[2] = 1;
|
|
} else {
|
|
globalWorkSize[0] = amd::alignUp(size[0], 8);
|
|
globalWorkSize[1] = amd::alignUp(size[1], 8);
|
|
globalWorkSize[2] = amd::alignUp(size[2], 4);
|
|
localWorkSize[0] = localWorkSize[1] = 8;
|
|
localWorkSize[2] = 4;
|
|
}
|
|
|
|
// Program kernels arguments for the blit operation
|
|
cl_mem clmem = ((cl_mem)as_cl<amd::Memory>(srcImage.owner()));
|
|
kernels_[blitType]->parameters().set(0, sizeof(cl_mem), &clmem);
|
|
clmem = ((cl_mem)as_cl<amd::Memory>(dstMemory.owner()));
|
|
kernels_[blitType]->parameters().set(1, sizeof(cl_mem), &clmem);
|
|
|
|
// Update extra paramters for USHORT and UBYTE pointers.
|
|
// Only then compiler can optimize the kernel to use
|
|
// UAV Raw for other writes
|
|
kernels_[blitType]->parameters().set(2, sizeof(cl_mem), &clmem);
|
|
kernels_[blitType]->parameters().set(3, sizeof(cl_mem), &clmem);
|
|
|
|
cl_int srcOrg[4] = {(cl_int)srcOrigin[0], (cl_int)srcOrigin[1],
|
|
(cl_int)srcOrigin[2], 0};
|
|
cl_int copySize[4] = {(cl_int)size[0], (cl_int)size[1], (cl_int)size[2], 0};
|
|
|
|
kernels_[blitType]->parameters().set(4, sizeof(srcOrg), srcOrg);
|
|
|
|
const size_t elementSize =
|
|
srcImage.owner()->asImage()->getImageFormat().getElementSize();
|
|
const size_t numChannels =
|
|
srcImage.owner()->asImage()->getImageFormat().getNumChannels();
|
|
|
|
// 1 element granularity for writes by default
|
|
cl_int granularity = 1;
|
|
if (elementSize == 2) {
|
|
granularity = 2;
|
|
} else if (elementSize >= 4) {
|
|
granularity = 4;
|
|
}
|
|
CondLog(((dstOrigin[0] % granularity) != 0), "Unaligned offset in blit!");
|
|
cl_ulong dstOrg[4] = {dstOrigin[0] / granularity, dstOrigin[1], dstOrigin[2],
|
|
0};
|
|
kernels_[blitType]->parameters().set(5, sizeof(dstOrg), dstOrg);
|
|
kernels_[blitType]->parameters().set(6, sizeof(copySize), copySize);
|
|
|
|
// Program memory format
|
|
uint multiplier = elementSize / sizeof(uint32_t);
|
|
multiplier = (multiplier == 0) ? 1 : multiplier;
|
|
cl_uint format[4] = {(cl_uint)numChannels,
|
|
(cl_uint)(elementSize / numChannels), multiplier, 0};
|
|
kernels_[blitType]->parameters().set(7, sizeof(format), format);
|
|
|
|
// Program row and slice pitches
|
|
cl_ulong pitch[4] = {0};
|
|
CalcRowSlicePitches(pitch, copySize, rowPitch, slicePitch, srcImage);
|
|
kernels_[blitType]->parameters().set(8, sizeof(pitch), pitch);
|
|
|
|
// Create ND range object for the kernel's execution
|
|
amd::NDRangeContainer ndrange(dim, globalWorkOffset, globalWorkSize,
|
|
localWorkSize);
|
|
|
|
// Execute the blit
|
|
address parameters = kernels_[blitType]->parameters().capture(dev_);
|
|
bool result = gpu().submitKernelInternal(ndrange, *kernels_[blitType],
|
|
parameters, NULL);
|
|
|
|
gpu().releaseGpuMemoryFence();
|
|
|
|
kernels_[blitType]->parameters().release(const_cast<address>(parameters),
|
|
dev_);
|
|
|
|
if (useView) {
|
|
srcView->owner()->release();
|
|
}
|
|
|
|
return result;
|
|
}
|
|
|
|
bool KernelBlitManager::copyBufferToImage(device::Memory& srcMemory,
|
|
device::Memory& dstMemory,
|
|
const amd::Coord3D& srcOrigin,
|
|
const amd::Coord3D& dstOrigin,
|
|
const amd::Coord3D& size, bool entire,
|
|
size_t rowPitch,
|
|
size_t slicePitch) const {
|
|
if (srcMemory.isHostMemDirectAccess()) {
|
|
return HsaBlitManager::copyBufferToImage(srcMemory, dstMemory, srcOrigin,
|
|
dstOrigin, size, entire, rowPitch,
|
|
slicePitch);
|
|
}
|
|
|
|
amd::Image::Format oldFormat = dstMemory.owner()->asImage()->getImageFormat();
|
|
amd::Image::Format newFormat = filterFormat(oldFormat);
|
|
bool useView = false;
|
|
|
|
device::Memory* dstView = &dstMemory;
|
|
if (oldFormat != newFormat) {
|
|
dstView = createImageView(dstMemory, newFormat);
|
|
useView = true;
|
|
}
|
|
|
|
roc::Image& dstImage = static_cast<roc::Image&>(*dstView);
|
|
|
|
// Use a common blit type with three dimensions by default
|
|
uint blitType = BlitCopyBufferToImage;
|
|
size_t dim = 0;
|
|
size_t globalWorkOffset[3] = {0, 0, 0};
|
|
size_t globalWorkSize[3];
|
|
size_t localWorkSize[3];
|
|
|
|
// Program the kernels workload depending on the blit dimensions
|
|
const size_t imageDims = dstImage.owner()->asImage()->getDims();
|
|
dim = 3;
|
|
if (imageDims == 1) {
|
|
globalWorkSize[0] = amd::alignUp(size[0], 256);
|
|
globalWorkSize[1] = amd::alignUp(size[1], 1);
|
|
globalWorkSize[2] = amd::alignUp(size[2], 1);
|
|
localWorkSize[0] = 256;
|
|
localWorkSize[1] = localWorkSize[2] = 1;
|
|
} else if (imageDims == 2) {
|
|
globalWorkSize[0] = amd::alignUp(size[0], 16);
|
|
globalWorkSize[1] = amd::alignUp(size[1], 16);
|
|
globalWorkSize[2] = amd::alignUp(size[2], 1);
|
|
localWorkSize[0] = localWorkSize[1] = 16;
|
|
localWorkSize[2] = 1;
|
|
} else {
|
|
globalWorkSize[0] = amd::alignUp(size[0], 8);
|
|
globalWorkSize[1] = amd::alignUp(size[1], 8);
|
|
globalWorkSize[2] = amd::alignUp(size[2], 4);
|
|
localWorkSize[0] = localWorkSize[1] = 8;
|
|
localWorkSize[2] = 4;
|
|
}
|
|
|
|
// Program kernels arguments for the blit operation
|
|
cl_mem clmem = ((cl_mem)as_cl<amd::Memory>(srcMemory.owner()));
|
|
kernels_[blitType]->parameters().set(0, sizeof(cl_mem), &clmem);
|
|
clmem = ((cl_mem)as_cl<amd::Memory>(dstImage.owner()));
|
|
kernels_[blitType]->parameters().set(1, sizeof(cl_mem), &clmem);
|
|
|
|
const size_t elementSize =
|
|
dstImage.owner()->asImage()->getImageFormat().getElementSize();
|
|
const size_t numChannels =
|
|
dstImage.owner()->asImage()->getImageFormat().getNumChannels();
|
|
|
|
// 1 element granularity for writes by default
|
|
cl_int granularity = 1;
|
|
if (elementSize == 2) {
|
|
granularity = 2;
|
|
} else if (elementSize >= 4) {
|
|
granularity = 4;
|
|
}
|
|
CondLog(((srcOrigin[0] % granularity) != 0), "Unaligned offset in blit!");
|
|
cl_ulong srcOrg[4] = {srcOrigin[0] / granularity, srcOrigin[1], srcOrigin[2],
|
|
0};
|
|
kernels_[blitType]->parameters().set(2, sizeof(srcOrg), srcOrg);
|
|
|
|
cl_int dstOrg[4] = {(cl_int)dstOrigin[0], (cl_int)dstOrigin[1],
|
|
(cl_int)dstOrigin[2], 0};
|
|
cl_int copySize[4] = {(cl_int)size[0], (cl_int)size[1], (cl_int)size[2], 0};
|
|
|
|
kernels_[blitType]->parameters().set(3, sizeof(dstOrg), dstOrg);
|
|
kernels_[blitType]->parameters().set(4, sizeof(copySize), copySize);
|
|
|
|
// Program memory format
|
|
uint multiplier = elementSize / sizeof(uint32_t);
|
|
multiplier = (multiplier == 0) ? 1 : multiplier;
|
|
cl_uint format[4] = {(cl_uint)numChannels,
|
|
(cl_uint)(elementSize / numChannels), multiplier, 0};
|
|
kernels_[blitType]->parameters().set(5, sizeof(format), format);
|
|
|
|
// Program row and slice pitches
|
|
cl_ulong pitch[4] = {0};
|
|
CalcRowSlicePitches(pitch, copySize, rowPitch, slicePitch, dstImage);
|
|
kernels_[blitType]->parameters().set(6, sizeof(pitch), pitch);
|
|
|
|
// Create ND range object for the kernel's execution
|
|
amd::NDRangeContainer ndrange(dim, globalWorkOffset, globalWorkSize,
|
|
localWorkSize);
|
|
|
|
// Execute the blit
|
|
address parameters = kernels_[blitType]->parameters().capture(dev_);
|
|
bool result = gpu().submitKernelInternal(ndrange, *kernels_[blitType],
|
|
parameters, NULL);
|
|
|
|
gpu().releaseGpuMemoryFence();
|
|
|
|
kernels_[blitType]->parameters().release(const_cast<address>(parameters),
|
|
dev_);
|
|
|
|
if (useView) {
|
|
dstView->owner()->release();
|
|
}
|
|
|
|
return result;
|
|
}
|
|
|
|
bool
|
|
KernelBlitManager::copyImage(
|
|
device::Memory& srcMemory,
|
|
device::Memory& dstMemory,
|
|
const amd::Coord3D& srcOrigin,
|
|
const amd::Coord3D& dstOrigin,
|
|
const amd::Coord3D& size,
|
|
bool entire) const
|
|
{
|
|
return HsaBlitManager::copyImage(
|
|
srcMemory, dstMemory, srcOrigin, dstOrigin, size, entire);
|
|
}
|
|
|
|
bool
|
|
KernelBlitManager::fillBuffer(
|
|
device::Memory& memory,
|
|
const void* pattern,
|
|
size_t patternSize,
|
|
const amd::Coord3D& origin,
|
|
const amd::Coord3D& size,
|
|
bool entire
|
|
) const
|
|
{
|
|
if (setup_.disableFillBuffer_ || memory.isHostMemDirectAccess()) {
|
|
return HostBlitManager::fillBuffer(memory, pattern, patternSize, origin,
|
|
size, entire);
|
|
}
|
|
|
|
uint fillType = FillBuffer;
|
|
size_t globalWorkOffset[3] = { 0, 0, 0 };
|
|
cl_ulong fillSize = size[0] / patternSize;
|
|
size_t globalWorkSize = amd::alignUp(fillSize, 256);
|
|
size_t localWorkSize = 256;
|
|
bool dwordAligned =
|
|
((patternSize % sizeof(uint32_t)) == 0) ? true : false;
|
|
|
|
// Program kernels arguments for the fill operation
|
|
if (dwordAligned) {
|
|
kernels_[fillType]->parameters().set(0, sizeof(cl_mem), NULL);
|
|
cl_mem clmem = ((cl_mem) as_cl<amd::Memory>(memory.owner()));
|
|
kernels_[fillType]->parameters().set(1, sizeof(cl_mem), &clmem);
|
|
}
|
|
else {
|
|
cl_mem clmem = ((cl_mem) as_cl<amd::Memory>(memory.owner()));
|
|
kernels_[fillType]->parameters().set(0, sizeof(cl_mem), &clmem);
|
|
kernels_[fillType]->parameters().set(1, sizeof(cl_mem), NULL);
|
|
}
|
|
|
|
amd::Buffer *fillMemory =
|
|
new (*context_) amd::Buffer(*context_, CL_MEM_USE_HOST_PTR, patternSize);
|
|
|
|
if (!fillMemory->create(const_cast<void *>(pattern))) {
|
|
LogError("[OCL] Fail to create mem object for destination");
|
|
return false;
|
|
}
|
|
|
|
if (fillMemory->getDeviceMemory(dev_) == NULL) {
|
|
LogError("[OCL] Fail to create device mem object for destination");
|
|
return false;
|
|
}
|
|
|
|
cl_mem clmem = ((cl_mem) as_cl<amd::Memory>(fillMemory));
|
|
kernels_[fillType]->parameters().set(2, sizeof(cl_mem), &clmem);
|
|
cl_ulong offset = origin[0];
|
|
if (dwordAligned) {
|
|
patternSize /= sizeof(uint32_t);
|
|
offset /= sizeof(uint32_t);
|
|
}
|
|
kernels_[fillType]->parameters().set(3, sizeof(cl_uint), &patternSize);
|
|
kernels_[fillType]->parameters().set(4, sizeof(offset), &offset);
|
|
kernels_[fillType]->parameters().set(5, sizeof(fillSize), &fillSize);
|
|
|
|
// Create ND range object for the kernel's execution
|
|
amd::NDRangeContainer ndrange(1,
|
|
globalWorkOffset, &globalWorkSize, &localWorkSize);
|
|
|
|
// Execute the blit
|
|
address parameters = kernels_[fillType]->parameters().capture(dev_);
|
|
bool result = gpu().submitKernelInternal(
|
|
ndrange, *kernels_[fillType], parameters, NULL);
|
|
kernels_[fillType]->parameters().release(const_cast<address>(parameters), dev_);
|
|
|
|
// Wait for the transfer to finish so that we could safely release the
|
|
// fill memory object.
|
|
// TODO: we could remove this if issue on implicit memory registration is
|
|
// fixed by KFD, so that we could pass the pattern as SVM.
|
|
gpu().releaseGpuMemoryFence();
|
|
|
|
fillMemory->release();
|
|
|
|
return result;
|
|
}
|
|
|
|
bool
|
|
KernelBlitManager::fillImage(
|
|
device::Memory& memory,
|
|
const void* pattern,
|
|
const amd::Coord3D& origin,
|
|
const amd::Coord3D& size,
|
|
bool entire
|
|
) const
|
|
{
|
|
return HsaBlitManager::fillImage(memory, pattern, origin, size, entire);
|
|
}
|
|
|
|
bool
|
|
KernelBlitManager::create(amd::Device& device)
|
|
{
|
|
if (!HsaBlitManager::create(device)) {
|
|
return false;
|
|
}
|
|
if (!createProgram(static_cast<Device&>(device))) {
|
|
return false;
|
|
}
|
|
|
|
return true;
|
|
}
|
|
|
|
bool
|
|
KernelBlitManager::createProgram(Device& device)
|
|
{
|
|
// Save context and program for this device
|
|
context_ = device.blitProgram()->context_;
|
|
context_->retain();
|
|
program_ = device.blitProgram()->program_;
|
|
program_->retain();
|
|
|
|
bool result = true;
|
|
|
|
// Create kernel objects for all blits
|
|
for (uint i = 0; i < BlitTotal; ++i) {
|
|
const amd::Symbol* symbol = program_->findSymbol(BlitName[i]);
|
|
if (symbol == NULL) {
|
|
result = false;
|
|
continue;
|
|
}
|
|
kernels_[i] = new amd::Kernel(*program_, *symbol, BlitName[i]);
|
|
if (kernels_[i] == NULL) {
|
|
result = false;
|
|
continue;
|
|
}
|
|
}
|
|
|
|
return result;
|
|
}
|
|
|
|
} // namespace roc
|