AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-432174 - Change the fillBuffer kernel

Browse Source 
- Add the new fillBuffer kernel, which allows to launch a limited
number of workgroups for memory fill operation
- Switch fill memory to 16 bytes write by default
- Allow to limit the workgroups with DEBUG_CLR_LIMIT_BLIT_WG

Change-Id: Ibad1822f2d42b2fc71bcfc1917c31409c0623e8e


[ROCm/clr commit: f1dc81f427]
This commit is contained in:
German Andryeyev
2023-11-14 12:49:17 -05:00
parent e1889b77b4
commit e390ec044f
13 changed files with 283 additions and 214 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/rocclr/device/blit.cpp
+41 -67
View File
@@ -344,7 +344,7 @@ bool HostBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory&
bool HostBlitManager::copyImageToBuffer(device::Memory& srcMemory, device::Memory& dstMemory, bool HostBlitManager::copyImageToBuffer(device::Memory& srcMemory, device::Memory& dstMemory,
const amd::Coord3D& srcOrigin, const amd::Coord3D& srcOrigin,
const amd::Coord3D& dstOrigin, const amd::Coord3D& size, const amd::Coord3D& dstOrigin, const amd::Coord3D& size,
bool entire, size_t rowPitch, size_t slicePitch, bool entire, size_t rowPitch, size_t slicePitch,
amd::CopyMetadata copyMetadata) const { amd::CopyMetadata copyMetadata) const {
size_t startLayer = srcOrigin[2]; size_t startLayer = srcOrigin[2];
size_t numLayers = size[2]; size_t numLayers = size[2];
@@ -475,7 +475,7 @@ bool HostBlitManager::copyBufferToImage(device::Memory& srcMemory, device::Memor
bool HostBlitManager::copyImage(device::Memory& srcMemory, device::Memory& dstMemory, bool HostBlitManager::copyImage(device::Memory& srcMemory, device::Memory& dstMemory,
const amd::Coord3D& srcOrigin, const amd::Coord3D& dstOrigin, const amd::Coord3D& srcOrigin, const amd::Coord3D& dstOrigin,
const amd::Coord3D& size, bool entire, const amd::Coord3D& size, bool entire,
amd::CopyMetadata copyMetadata) const { amd::CopyMetadata copyMetadata) const {
size_t startLayer = srcOrigin[2]; size_t startLayer = srcOrigin[2];
size_t numLayers = size[2]; size_t numLayers = size[2];
@@ -685,49 +685,41 @@ uint32_t HostBlitManager::sRGBmap(float fc) const {
return (uint32_t)(c * 255.0 + 0.5); return (uint32_t)(c * 255.0 + 0.5);
} }
bool HostBlitManager::FillBufferInfo::ExpandPattern64(uint64_t pattern, size_t pattern_size, // ================================================================================================
uint64_t& pattern64) { void HostBlitManager::FillBufferInfo::ExpandPattern(uint32_t pattern_size, const void* pattern) {
// If pattern size exceeds extended, then runtime will select the normal path
if (pattern_size >= kExtendedSize) {
return;
}
bool retval = true; pattern_expanded_ = true;
if (pattern_size == sizeof(uint8_t)) {
do { uint8_t pattern_byte = *reinterpret_cast<const uint8_t*>(pattern);
for (uint32_t i = 0; i < kExtendedSize; ++i) {
// If the pattern is 0 or if the pattern_size is same as max size. reinterpret_cast<uint8_t*>(expanded_pattern_)[i] = pattern_byte;
if (pattern == 0 || pattern_size == sizeof(uint64_t)) {
pattern64 = pattern;
break;
} }
} else if (pattern_size == sizeof(uint16_t)) {
// Clean Curr_pattern, since it was casted off from const void* with a lesser size than size_t. uint16_t pattern_word = *reinterpret_cast<const uint16_t*>(pattern);
ClearBits64(pattern, (pattern_size * 8)); for (uint32_t i = 0; i < kExtendedSize / sizeof(uint16_t); ++i) {
pattern64 = 0; reinterpret_cast<uint16_t*>(expanded_pattern_)[i] = pattern_word;
if (pattern_size == sizeof(uint8_t)) {
pattern = pattern & 0xff;
pattern64 = ((pattern << 56) | (pattern << 48) | (pattern << 40) | (pattern << 32)
| (pattern << 24) | (pattern << 16) | (pattern << 8) | (pattern));
} else if (pattern_size == sizeof(uint16_t)) {
pattern = pattern & 0xffff;
pattern64 = ((pattern << 48) | (pattern << 32) | (pattern << 16) | (pattern));
} else if (pattern_size == sizeof(uint32_t)) {
pattern = pattern & 0xffffffff;
pattern64 = ((pattern << 32) | (pattern));
} else {
LogPrintfError("Unsupported Pattern size: %u \n", pattern_size);
retval = false;
break;
} }
} else if (pattern_size == sizeof(uint32_t)) {
} while (0); uint32_t pattern_dword = *reinterpret_cast<const uint32_t*>(pattern);
for (uint32_t i = 0; i < kExtendedSize / sizeof(uint32_t); ++i) {
return retval; reinterpret_cast<uint32_t*>(expanded_pattern_)[i] = pattern_dword;
}
} else {
uint64_t pattern_qword = *reinterpret_cast<const uint64_t*>(pattern);
reinterpret_cast<uint64_t*>(expanded_pattern_)[0] = pattern_qword;
reinterpret_cast<uint64_t*>(expanded_pattern_)[1] = pattern_qword;
}
} }
bool HostBlitManager::FillBufferInfo::PackInfo(const device::Memory& memory, size_t fill_size, // ================================================================================================
size_t fill_origin, const void* pattern_ptr, void HostBlitManager::FillBufferInfo::PackInfo(const device::Memory& memory, size_t fill_size,
size_t pattern_size, size_t fill_origin, const void* pattern_ptr,
std::vector<FillBufferInfo>& packed_info) { size_t pattern_size,
std::vector<FillBufferInfo>& packed_info) {
// 1. Validate input arguments // 1. Validate input arguments
guarantee(fill_size >= pattern_size, "Pattern Size: %u cannot be greater than fill size: %u \n", guarantee(fill_size >= pattern_size, "Pattern Size: %u cannot be greater than fill size: %u \n",
pattern_size, fill_size); pattern_size, fill_size);
@@ -736,60 +728,42 @@ bool HostBlitManager::FillBufferInfo::PackInfo(const device::Memory& memory, siz
// 2. Calculate the next closest dword aligned address for faster processing // 2. Calculate the next closest dword aligned address for faster processing
size_t dst_addr = memory.virtualAddress() + fill_origin; size_t dst_addr = memory.virtualAddress() + fill_origin;
size_t aligned_dst_addr = amd::alignUp(dst_addr, sizeof(size_t)); size_t aligned_dst_addr = amd::alignUp(dst_addr, kExtendedSize);
guarantee(aligned_dst_addr >= dst_addr, "Aligned address: %u cannot be greater than destination" guarantee(aligned_dst_addr >= dst_addr, "Aligned address: %u cannot be greater than destination"
"address :%u \n", aligned_dst_addr, dst_addr); "address :%u \n", aligned_dst_addr, dst_addr);
// 3. If given address is not aligned calculate head and tail size. // 3. If given address is not aligned calculate head and tail size.
size_t head_size = std::min(aligned_dst_addr - dst_addr, fill_size); size_t head_size = std::min(aligned_dst_addr - dst_addr, fill_size);
size_t aligned_size = ((fill_size - head_size) / sizeof(size_t)) * sizeof(size_t); size_t aligned_size = ((fill_size - head_size) / kExtendedSize) * kExtendedSize;
size_t tail_size = (fill_size - head_size) % sizeof(size_t); size_t tail_size = (fill_size - head_size) % kExtendedSize;
guarantee((head_size + aligned_size + tail_size) <= fill_size, "Head size, aligned size & tail" guarantee((head_size + aligned_size + tail_size) <= fill_size, "Head size, aligned size & tail"
"size together cannot cross fill size"); "size together cannot cross fill size");
// 4. Clear unwanted bytes from the pattern if the pattern size is < sizeof(size_t). // 4. Fill the head, aligned, tail info if they exist.
uint64_t pattern = *(reinterpret_cast<uint64_t*>(const_cast<void*>(pattern_ptr)));
if (pattern_size < sizeof(uint64_t)) {
ClearBits64(pattern, (pattern_size * 8));
}
// 5. Fill the head, aligned, tail info if they exist.
FillBufferInfo fill_info;
if (head_size > 0) { if (head_size > 0) {
// Offsetted ptrs should align with pattern size. Runtime not responsible for rotating pattern. // Offsetted ptrs should align with pattern size. Runtime not responsible for rotating pattern.
guarantee((head_size % pattern_size) == 0, "Offseted ptr should align with pattern_size"); guarantee((head_size % pattern_size) == 0, "Offseted ptr should align with pattern_size");
fill_info.fill_size_ = head_size; FillBufferInfo fill_info(head_size);
packed_info.push_back(fill_info); packed_info.push_back(fill_info);
} }
fill_info.clearInfo();
if (aligned_size > 0) { if (aligned_size > 0) {
// Offsetted ptrs should align with pattern size. Runtime not responsible for rotating pattern. // Offsetted ptrs should align with pattern size. Runtime not responsible for rotating pattern.
guarantee((aligned_size % pattern_size) == 0, "Offseted ptr should align with pattern_size"); guarantee((aligned_size % pattern_size) == 0, "Offseted ptr should align with pattern_size");
if (pattern_size < sizeof(uint64_t)) { FillBufferInfo fill_info(aligned_size);
if (!ExpandPattern64(pattern, pattern_size, fill_info.expanded_pattern_)) { fill_info.ExpandPattern(pattern_size, pattern_ptr);
DevLogPrintfError("Failed Expanding the pattern for pattern:%u, pattern_size: %u",
pattern, pattern_size);
return false;
}
fill_info.pattern_expanded_ = true;
}
fill_info.fill_size_ = aligned_size;
packed_info.push_back(fill_info); packed_info.push_back(fill_info);
} }
fill_info.clearInfo();
if (tail_size > 0) { if (tail_size > 0) {
// Offsetted ptrs should align with pattern size. Runtime not responsible for rotating pattern. // Offsetted ptrs should align with pattern size. Runtime not responsible for rotating pattern.
guarantee((tail_size % pattern_size) == 0, "Offseted ptr should align with pattern_size"); guarantee((tail_size % pattern_size) == 0, "Offseted ptr should align with pattern_size");
fill_info.fill_size_ = tail_size; FillBufferInfo fill_info(tail_size);
packed_info.push_back(fill_info); packed_info.push_back(fill_info);
} }
fill_info.clearInfo();
return true;
} }
} // namespace gpu } // namespace gpu
projects/clr/rocclr/device/blit.hpp
+21 -33
View File
@@ -78,7 +78,7 @@ class BlitManager : public amd::HeapObject {
const amd::Coord3D& origin, //!< Source origin const amd::Coord3D& origin, //!< Source origin
const amd::Coord3D& size, //!< Size of the copy region const amd::Coord3D& size, //!< Size of the copy region
bool entire = false, //!< Entire buffer will be updated bool entire = false, //!< Entire buffer will be updated
amd::CopyMetadata copyMetadata = amd::CopyMetadata copyMetadata =
amd::CopyMetadata()//!< Memory copy MetaData amd::CopyMetadata()//!< Memory copy MetaData
) const = 0; ) const = 0;
@@ -89,7 +89,7 @@ class BlitManager : public amd::HeapObject {
const amd::BufferRect& hostRect, //!< Destination rectangle const amd::BufferRect& hostRect, //!< Destination rectangle
const amd::Coord3D& size, //!< Size of the copy region const amd::Coord3D& size, //!< Size of the copy region
bool entire = false, //!< Entire buffer will be updated bool entire = false, //!< Entire buffer will be updated
amd::CopyMetadata copyMetadata = amd::CopyMetadata copyMetadata =
amd::CopyMetadata() //!< Memory copy MetaData amd::CopyMetadata() //!< Memory copy MetaData
) const = 0; ) const = 0;
@@ -101,7 +101,7 @@ class BlitManager : public amd::HeapObject {
size_t rowPitch, //!< Row pitch for host memory size_t rowPitch, //!< Row pitch for host memory
size_t slicePitch, //!< Slice pitch for host memory size_t slicePitch, //!< Slice pitch for host memory
bool entire = false, //!< Entire buffer will be updated bool entire = false, //!< Entire buffer will be updated
amd::CopyMetadata copyMetadata = amd::CopyMetadata copyMetadata =
amd::CopyMetadata()//!< Memory copy MetaData amd::CopyMetadata()//!< Memory copy MetaData
) const = 0; ) const = 0;
@@ -111,7 +111,7 @@ class BlitManager : public amd::HeapObject {
const amd::Coord3D& origin, //!< Destination origin const amd::Coord3D& origin, //!< Destination origin
const amd::Coord3D& size, //!< Size of the copy region const amd::Coord3D& size, //!< Size of the copy region
bool entire = false, //!< Entire buffer will be updated bool entire = false, //!< Entire buffer will be updated
amd::CopyMetadata copyMetadata = amd::CopyMetadata copyMetadata =
amd::CopyMetadata() //!< Memory copy MetaData amd::CopyMetadata() //!< Memory copy MetaData
) const = 0; ) const = 0;
@@ -122,7 +122,7 @@ class BlitManager : public amd::HeapObject {
const amd::BufferRect& bufRect, //!< Source rectangle const amd::BufferRect& bufRect, //!< Source rectangle
const amd::Coord3D& size, //!< Size of the copy region const amd::Coord3D& size, //!< Size of the copy region
bool entire = false, //!< Entire buffer will be updated bool entire = false, //!< Entire buffer will be updated
amd::CopyMetadata copyMetadata = amd::CopyMetadata copyMetadata =
amd::CopyMetadata() //!< Memory copy MetaData amd::CopyMetadata() //!< Memory copy MetaData
) const = 0; ) const = 0;
@@ -134,7 +134,7 @@ class BlitManager : public amd::HeapObject {
size_t rowPitch, //!< Row pitch for host memory size_t rowPitch, //!< Row pitch for host memory
size_t slicePitch, //!< Slice pitch for host memory size_t slicePitch, //!< Slice pitch for host memory
bool entire = false, //!< Entire buffer will be updated bool entire = false, //!< Entire buffer will be updated
amd::CopyMetadata copyMetadata = amd::CopyMetadata copyMetadata =
amd::CopyMetadata() //!< Memory copy MetaData amd::CopyMetadata() //!< Memory copy MetaData
) const = 0; ) const = 0;
@@ -193,7 +193,7 @@ class BlitManager : public amd::HeapObject {
const amd::Coord3D& dstOrigin, //!< Destination origin const amd::Coord3D& dstOrigin, //!< Destination origin
const amd::Coord3D& size, //!< Size of the copy region const amd::Coord3D& size, //!< Size of the copy region
bool entire = false, //!< Entire buffer will be updated bool entire = false, //!< Entire buffer will be updated
amd::CopyMetadata copyMetadata = amd::CopyMetadata copyMetadata =
amd::CopyMetadata() //!< Memory copy MetaData amd::CopyMetadata() //!< Memory copy MetaData
) const = 0; ) const = 0;
@@ -301,7 +301,7 @@ class HostBlitManager : public device::BlitManager {
size_t rowPitch, //!< Row pitch for host memory size_t rowPitch, //!< Row pitch for host memory
size_t slicePitch, //!< Slice pitch for host memory size_t slicePitch, //!< Slice pitch for host memory
bool entire = false, //!< Entire buffer will be updated bool entire = false, //!< Entire buffer will be updated
amd::CopyMetadata copyMetadata = amd::CopyMetadata copyMetadata =
amd::CopyMetadata() //!< Memory copy MetaData amd::CopyMetadata() //!< Memory copy MetaData
) const; ) const;
@@ -334,7 +334,7 @@ class HostBlitManager : public device::BlitManager {
size_t rowPitch, //!< Row pitch for host memory size_t rowPitch, //!< Row pitch for host memory
size_t slicePitch, //!< Slice pitch for host memory size_t slicePitch, //!< Slice pitch for host memory
bool entire = false, //!< Entire buffer will be updated bool entire = false, //!< Entire buffer will be updated
amd::CopyMetadata copyMetadata = amd::CopyMetadata copyMetadata =
amd::CopyMetadata() //!< Memory copy MetaData amd::CopyMetadata() //!< Memory copy MetaData
) const; ) const;
@@ -393,7 +393,7 @@ class HostBlitManager : public device::BlitManager {
const amd::Coord3D& dstOrigin, //!< Destination origin const amd::Coord3D& dstOrigin, //!< Destination origin
const amd::Coord3D& size, //!< Size of the copy region const amd::Coord3D& size, //!< Size of the copy region
bool entire = false, //!< Entire buffer will be updated bool entire = false, //!< Entire buffer will be updated
amd::CopyMetadata copyMetadata = amd::CopyMetadata copyMetadata =
amd::CopyMetadata() //!< Memory copy MetaData amd::CopyMetadata() //!< Memory copy MetaData
) const; ) const;
@@ -423,36 +423,24 @@ class HostBlitManager : public device::BlitManager {
const amd::Device& dev_; //!< Physical device const amd::Device& dev_; //!< Physical device
// Packed Fill Buffer // Packed Fill Buffer
class FillBufferInfo { struct FillBufferInfo {
public: static constexpr uint32_t kExtendedSize = 2 * sizeof(uint64_t);
FillBufferInfo(): fill_size_(0), expanded_pattern_(0), pattern_expanded_(false) {}
static bool PackInfo(const device::Memory& memory, size_t fill_size, static void PackInfo(const device::Memory& memory, size_t fill_size,
size_t fill_origin, const void* pattern, size_t pattern_size, size_t fill_origin, const void* pattern, size_t pattern_size,
std::vector<FillBufferInfo>& packed_info); std::vector<FillBufferInfo>& packed_info);
private:
static bool ExpandPattern64(uint64_t pattern, size_t pattern_size, uint64_t& pattern64);
static inline void ClearBits64(uint64_t& pattern, uint64_t num_bits) {
pattern &= ~(~(static_cast<uint64_t>(0)) << num_bits);
FillBufferInfo(size_t fill_size): fill_size_(fill_size), pattern_expanded_(false) {
memset(&expanded_pattern_, 0, sizeof(expanded_pattern_));
} }
void clearInfo () { void ExpandPattern(uint32_t pattern_size, const void* pattern);
fill_size_ = 0;
expanded_pattern_ = 0;
pattern_expanded_ = false;
}
public: size_t fill_size_; //!< Fill size for this command
size_t fill_size_; // Fill size for this command uint8_t expanded_pattern_[kExtendedSize]; //!< Pattern for this command - 16 bytes
uint64_t expanded_pattern_; // Pattern for this command bool pattern_expanded_; //!< Boolean to check if pattern is expanded
bool pattern_expanded_; // Boolean to check if pattern is expanded
}; };
private: private:
//! Disable copy constructor //! Disable copy constructor
HostBlitManager(const HostBlitManager&); HostBlitManager(const HostBlitManager&);
projects/clr/rocclr/device/blitcl.cpp
+51 -10
View File
@@ -47,16 +47,57 @@ const char* BlitLinearSourceCode = BLIT_KERNELS(
extern void __ockl_gws_init(uint nwm1, uint rid); extern void __ockl_gws_init(uint nwm1, uint rid);
// Implementation __kernel void __amd_rocclr_fillBufferAligned(
__kernel void __amd_rocclr_fillBufferAligned(__global uchar* bufUChar, __global uchar* bufUChar, __global ushort* bufUShort, __global uint* bufUInt,
__global ushort* bufUShort, __global ulong* bufULong, __global ulong2* bufULong2, __constant uchar* pattern,
__global uint* bufUInt, uint pattern_size, ulong offset, ulong end_ptr, uint next_chunk) {
__global ulong* bufULong, int id = get_global_id(0);
__constant uchar* pattern, long cur_id = offset + id * pattern_size;
uint patternSize, ulong offset, if (bufULong2) {
ulong size) { __global ulong2* element = &bufULong2[cur_id];
__amd_fillBufferAligned(bufUChar, bufUShort, bufUInt, bufULong, __constant ulong2* pt = (__constant ulong2*)pattern;
pattern, patternSize, offset, size); while ((ulong)element < end_ptr) {
for (uint i = 0; i < pattern_size; ++i) {
element[i] = pt[i];
}
element += next_chunk;
}
} else if (bufULong) {
__global ulong* element = &bufULong[cur_id];
__constant ulong* pt = (__constant ulong*)pattern;
while ((ulong)element < end_ptr) {
for (uint i = 0; i < pattern_size; ++i) {
element[i] = pt[i];
}
element += next_chunk;
}
} else if (bufUInt) {
__global uint* element = &bufUInt[cur_id];
__constant uint* pt = (__constant uint*)pattern;
while ((ulong)element < end_ptr) {
for (uint i = 0; i < pattern_size; ++i) {
element[i] = pt[i];
}
element += next_chunk;
}
} else if (bufUShort) {
__global ushort* element = &bufUShort[cur_id];
__constant ushort* pt = (__constant ushort*)pattern;
while ((ulong)element < end_ptr) {
for (uint i = 0; i < pattern_size; ++i) {
element[i] = pt[i];
}
element += next_chunk;
}
} else {
__global uchar* element = &bufUChar[cur_id];
while ((ulong)element < end_ptr) {
for (uint i = 0; i < pattern_size; ++i) {
element[i] = pattern[i];
}
element += next_chunk;
}
}
} }
__kernel void __amd_rocclr_fillBufferAligned2D(__global uchar* bufUChar, __kernel void __amd_rocclr_fillBufferAligned2D(__global uchar* bufUChar,
projects/clr/rocclr/device/pal/palblit.cpp
+88 -55
View File
@@ -2172,63 +2172,96 @@ bool KernelBlitManager::fillBuffer(device::Memory& memory, const void* pattern,
synchronize(); synchronize();
return result; return result;
} else { } else {
uint fillType = FillBufferAligned; // Pack the fill buffer info, that handles unaligned memories.
size_t globalWorkOffset[3] = {0, 0, 0}; std::vector<FillBufferInfo> packed_vector{};
uint64_t fillSize = size[0] / patternSize; FillBufferInfo::PackInfo(memory, size[0], origin[0], pattern, patternSize, packed_vector);
size_t globalWorkSize = amd::alignUp(fillSize, 256);
size_t localWorkSize = 256;
uint32_t alignment = (patternSize & 0x7) == 0 ?
sizeof(uint64_t) :
(patternSize & 0x3) == 0 ?
sizeof(uint32_t) :
(patternSize & 0x1) == 0 ?
sizeof(uint16_t) : sizeof(uint8_t);
// Program kernels arguments for the fill operation size_t overall_offset = origin[0];
Memory* mem = &gpuMem(memory); for (auto& packed_obj : packed_vector) {
if (alignment == sizeof(uint64_t)) { constexpr uint32_t kFillType = FillBufferAligned;
setArgument(kernels_[fillType], 0, sizeof(cl_mem), nullptr); uint32_t kpattern_size = (packed_obj.pattern_expanded_)
setArgument(kernels_[fillType], 1, sizeof(cl_mem), nullptr); ? HostBlitManager::FillBufferInfo::kExtendedSize
setArgument(kernels_[fillType], 2, sizeof(cl_mem), nullptr); : patternSize;
setArgument(kernels_[fillType], 3, sizeof(cl_mem), &mem); size_t kfill_size = packed_obj.fill_size_ / kpattern_size;
} else if (alignment == sizeof(uint32_t)) { size_t koffset = overall_offset;
setArgument(kernels_[fillType], 0, sizeof(cl_mem), nullptr); overall_offset += packed_obj.fill_size_;
setArgument(kernels_[fillType], 1, sizeof(cl_mem), nullptr);
setArgument(kernels_[fillType], 2, sizeof(cl_mem), &mem); size_t globalWorkOffset[3] = {0, 0, 0};
setArgument(kernels_[fillType], 3, sizeof(cl_mem), nullptr); uint32_t alignment = (kpattern_size & 0xf) == 0 ? 2 * sizeof(uint64_t) :
} else if (alignment == sizeof(uint16_t)) { (kpattern_size & 0x7) == 0 ? sizeof(uint64_t) :
setArgument(kernels_[fillType], 0, sizeof(cl_mem), nullptr); (kpattern_size & 0x3) == 0 ? sizeof(uint32_t) :
setArgument(kernels_[fillType], 1, sizeof(cl_mem), &mem); (kpattern_size & 0x1) == 0 ? sizeof(uint16_t) : sizeof(uint8_t);
setArgument(kernels_[fillType], 2, sizeof(cl_mem), nullptr);
setArgument(kernels_[fillType], 3, sizeof(cl_mem), nullptr); // Program kernels arguments for the fill operation
} else { Memory* mem = &gpuMem(memory);
setArgument(kernels_[fillType], 0, sizeof(cl_mem), &mem); if (alignment == 2 * sizeof(uint64_t)) {
setArgument(kernels_[fillType], 1, sizeof(cl_mem), nullptr); setArgument(kernels_[kFillType], 0, sizeof(cl_mem), nullptr);
setArgument(kernels_[fillType], 2, sizeof(cl_mem), nullptr); setArgument(kernels_[kFillType], 1, sizeof(cl_mem), nullptr);
setArgument(kernels_[fillType], 3, sizeof(cl_mem), nullptr); setArgument(kernels_[kFillType], 2, sizeof(cl_mem), nullptr);
setArgument(kernels_[kFillType], 3, sizeof(cl_mem), nullptr);
setArgument(kernels_[kFillType], 4, sizeof(cl_mem), &mem);
} else if (alignment == sizeof(uint64_t)) {
setArgument(kernels_[kFillType], 0, sizeof(cl_mem), nullptr);
setArgument(kernels_[kFillType], 1, sizeof(cl_mem), nullptr);
setArgument(kernels_[kFillType], 2, sizeof(cl_mem), nullptr);
setArgument(kernels_[kFillType], 3, sizeof(cl_mem), &mem);
setArgument(kernels_[kFillType], 4, sizeof(cl_mem), nullptr);
} else if (alignment == sizeof(uint32_t)) {
setArgument(kernels_[kFillType], 0, sizeof(cl_mem), nullptr);
setArgument(kernels_[kFillType], 1, sizeof(cl_mem), nullptr);
setArgument(kernels_[kFillType], 2, sizeof(cl_mem), &mem);
setArgument(kernels_[kFillType], 3, sizeof(cl_mem), nullptr);
setArgument(kernels_[kFillType], 4, sizeof(cl_mem), nullptr);
} else if (alignment == sizeof(uint16_t)) {
setArgument(kernels_[kFillType], 0, sizeof(cl_mem), nullptr);
setArgument(kernels_[kFillType], 1, sizeof(cl_mem), &mem);
setArgument(kernels_[kFillType], 2, sizeof(cl_mem), nullptr);
setArgument(kernels_[kFillType], 3, sizeof(cl_mem), nullptr);
setArgument(kernels_[kFillType], 4, sizeof(cl_mem), nullptr);
} else {
setArgument(kernels_[kFillType], 0, sizeof(cl_mem), &mem);
setArgument(kernels_[kFillType], 1, sizeof(cl_mem), nullptr);
setArgument(kernels_[kFillType], 2, sizeof(cl_mem), nullptr);
setArgument(kernels_[kFillType], 3, sizeof(cl_mem), nullptr);
setArgument(kernels_[kFillType], 4, sizeof(cl_mem), nullptr);
}
const size_t localWorkSize = 256;
size_t globalWorkSize =
std::min(dev().settings().limit_blit_wg_ * localWorkSize, kfill_size);
globalWorkSize = amd::alignUp(globalWorkSize, localWorkSize);
Memory& gpuCB = gpu().xferWrite().Acquire(patternSize);
void* constBuf = gpuCB.map(&gpu(), Resource::NoWait);
// If pattern has been expanded, use the expanded pattern, otherwise use the default pattern
if (packed_obj.pattern_expanded_) {
memcpy(constBuf, &packed_obj.expanded_pattern_, kpattern_size);
} else {
memcpy(constBuf, pattern, kpattern_size);
}
gpuCB.unmap(&gpu());
Memory* pGpuCB = &gpuCB;
setArgument(kernels_[kFillType], 5, sizeof(cl_mem), &pGpuCB);
uint64_t offset = origin[0];
// Adjust the pattern size in the copy type size
kpattern_size /= alignment;
setArgument(kernels_[kFillType], 6, sizeof(uint32_t), &kpattern_size);
koffset /= alignment;
setArgument(kernels_[kFillType], 7, sizeof(koffset), &koffset);
// Calculate max id
kfill_size = memory.virtualAddress() + (koffset + kfill_size * kpattern_size) * alignment;
setArgument(kernels_[kFillType], 8, sizeof(kfill_size), &kfill_size);
uint32_t next_chunk = globalWorkSize * kpattern_size;
setArgument(kernels_[kFillType], 9, sizeof(uint32_t), &next_chunk);
// Create ND range object for the kernel's execution
amd::NDRangeContainer ndrange(1, globalWorkOffset, &globalWorkSize, &localWorkSize);
// Execute the blit
address parameters = kernels_[kFillType]->parameters().values();
result = gpu().submitKernelInternal(ndrange, *kernels_[kFillType], parameters);
gpu().xferWrite().Release(gpuCB);
} }
Memory& gpuCB = gpu().xferWrite().Acquire(patternSize);
void* constBuf = gpuCB.map(&gpu(), Resource::NoWait);
memcpy(constBuf, pattern, patternSize);
gpuCB.unmap(&gpu());
Memory* pGpuCB = &gpuCB;
setArgument(kernels_[fillType], 4, sizeof(cl_mem), &pGpuCB);
uint64_t offset = origin[0];
patternSize/= alignment;
offset /= alignment;
setArgument(kernels_[fillType], 5, sizeof(uint32_t), &patternSize);
setArgument(kernels_[fillType], 6, sizeof(offset), &offset);
setArgument(kernels_[fillType], 7, 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().values();
result = gpu().submitKernelInternal(ndrange, *kernels_[fillType], parameters);
gpu().xferWrite().Release(gpuCB);
} }
synchronize(); synchronize();
projects/clr/rocclr/device/pal/palsettings.cpp
+11 -1
View File
@@ -148,6 +148,8 @@ Settings::Settings() {
prepinnedMinSize_ = 0; prepinnedMinSize_ = 0;
cpDmaCopySizeMax_ = GPU_CP_DMA_COPY_SIZE * Ki; cpDmaCopySizeMax_ = GPU_CP_DMA_COPY_SIZE * Ki;
useDeviceKernelArg_ = flagIsDefault(HIP_FORCE_DEV_KERNARG) ? false : HIP_FORCE_DEV_KERNARG; useDeviceKernelArg_ = flagIsDefault(HIP_FORCE_DEV_KERNARG) ? false : HIP_FORCE_DEV_KERNARG;
limit_blit_wg_ = 16;
} }
bool Settings::create(const Pal::DeviceProperties& palProp, bool Settings::create(const Pal::DeviceProperties& palProp,
@@ -288,7 +290,7 @@ bool Settings::create(const Pal::DeviceProperties& palProp,
// Cap at OpenCL20 for now // Cap at OpenCL20 for now
if (oclVersion_ > OpenCL20) oclVersion_ = OpenCL20; if (oclVersion_ > OpenCL20) oclVersion_ = OpenCL20;
use64BitPtr_ = LP64_SWITCH(false, true); use64BitPtr_ = LP64_SWITCH(false, true);
if (oclVersion_ >= OpenCL20) { if (oclVersion_ >= OpenCL20) {
@@ -435,6 +437,10 @@ bool Settings::create(const Pal::DeviceProperties& palProp,
prepinnedMinSize_ = PAL_PREPINNED_MEMORY_SIZE * Ki; prepinnedMinSize_ = PAL_PREPINNED_MEMORY_SIZE * Ki;
} }
limit_blit_wg_ = enableWgpMode_
? palProp.gfxipProperties.shaderCore.numAvailableCus / 2
: palProp.gfxipProperties.shaderCore.numAvailableCus;
// Override current device settings // Override current device settings
override(); override();
@@ -501,6 +507,10 @@ void Settings::override() {
if (!flagIsDefault(PAL_ALWAYS_RESIDENT)) { if (!flagIsDefault(PAL_ALWAYS_RESIDENT)) {
alwaysResident_ = PAL_ALWAYS_RESIDENT; alwaysResident_ = PAL_ALWAYS_RESIDENT;
} }
if (!flagIsDefault(DEBUG_CLR_LIMIT_BLIT_WG)) {
limit_blit_wg_ = std::max(DEBUG_CLR_LIMIT_BLIT_WG, 0x1U);
}
} }
} // namespace pal } // namespace pal
projects/clr/rocclr/device/pal/palsettings.hpp
+1
View File
@@ -118,6 +118,7 @@ class Settings : public device::Settings {
amd::LibrarySelector libSelector_; //!< Select linking libraries for compiler amd::LibrarySelector libSelector_; //!< Select linking libraries for compiler
size_t prepinnedMinSize_; //!< minimal memory size for prepinned transfer size_t prepinnedMinSize_; //!< minimal memory size for prepinned transfer
uint32_t limit_blit_wg_; //!< The number of workgroups for blit execution
//! Default constructor //! Default constructor
Settings(); Settings();
projects/clr/rocclr/device/pal/palvirtual.cpp
+3
View File
@@ -778,6 +778,9 @@ bool VirtualGPU::createVirtualQueue(uint deviceQueueSize) {
// Add mask array for AmdAqlWrap slots // Add mask array for AmdAqlWrap slots
allocSize += amd::alignUp(numSlots, DeviceQueueMaskSize) / 8; allocSize += amd::alignUp(numSlots, DeviceQueueMaskSize) / 8;
// Align size to 64 bytes for more efficient fill operation
allocSize = amd::alignUp(allocSize, 8 * sizeof(uint64_t));
virtualQueue_ = new Memory(dev(), allocSize); virtualQueue_ = new Memory(dev(), allocSize);
Resource::MemoryType type = (GPU_PRINT_CHILD_KERNEL == 0) ? Resource::Local : Resource::Remote; Resource::MemoryType type = (GPU_PRINT_CHILD_KERNEL == 0) ? Resource::Local : Resource::Remote;
if ((virtualQueue_ == nullptr) || !virtualQueue_->create(type)) { if ((virtualQueue_ == nullptr) || !virtualQueue_->create(type)) {
projects/clr/rocclr/device/rocm/rocblit.cpp
+57 -45
View File
@@ -2048,9 +2048,10 @@ bool KernelBlitManager::fillBuffer(device::Memory& memory, const void* pattern,
} }
// ================================================================================================ // ================================================================================================
bool KernelBlitManager::fillBuffer1D(device::Memory& memory, const void* pattern, size_t patternSize, bool KernelBlitManager::fillBuffer1D(device::Memory& memory, const void* pattern,
const amd::Coord3D& surface, const amd::Coord3D& origin, size_t patternSize, const amd::Coord3D& surface,
const amd::Coord3D& size, bool entire, bool forceBlit) const { const amd::Coord3D& origin, const amd::Coord3D& size,
bool entire, bool forceBlit) const {
amd::ScopedLock k(lockXferOps_); amd::ScopedLock k(lockXferOps_);
bool result = false; bool result = false;
@@ -2063,79 +2064,90 @@ bool KernelBlitManager::fillBuffer1D(device::Memory& memory, const void* pattern
synchronize(); synchronize();
return result; return result;
} else { } else {
// Pack the fill buffer info, that handles unaligned memories. // Pack the fill buffer info, that handles unaligned memories.
std::vector<FillBufferInfo> packed_vector{}; std::vector<FillBufferInfo> packed_vector{};
FillBufferInfo::PackInfo(memory, size[0], origin[0], pattern, patternSize, packed_vector); FillBufferInfo::PackInfo(memory, size[0], origin[0], pattern, patternSize, packed_vector);
size_t overall_offset = origin[0]; size_t overall_offset = origin[0];
for (auto& packed_obj: packed_vector) { for (auto& packed_obj: packed_vector) {
uint fillType = FillBufferAligned; constexpr uint32_t kFillType = FillBufferAligned;
uint32_t kpattern_size = (packed_obj.pattern_expanded_) ?
uint32_t kpattern_size32 = (packed_obj.pattern_expanded_) ? sizeof(size_t) : patternSize; HostBlitManager::FillBufferInfo::kExtendedSize : patternSize;
size_t kfill_size = packed_obj.fill_size_/kpattern_size32; size_t kfill_size = packed_obj.fill_size_ / kpattern_size;
size_t koffset = overall_offset; size_t koffset = overall_offset;
overall_offset += packed_obj.fill_size_; overall_offset += packed_obj.fill_size_;
size_t globalWorkOffset[3] = {0, 0, 0}; size_t globalWorkOffset[3] = {0, 0, 0};
size_t globalWorkSize = amd::alignUp(kfill_size, 256); uint32_t alignment = (kpattern_size & 0xf) == 0 ? 2 * sizeof(uint64_t) :
size_t localWorkSize = 256; (kpattern_size & 0x7) == 0 ? sizeof(uint64_t) :
(kpattern_size & 0x3) == 0 ? sizeof(uint32_t) :
uint32_t alignment = (kpattern_size32 & 0x7) == 0 ? (kpattern_size & 0x1) == 0 ? sizeof(uint16_t) : sizeof(uint8_t);
sizeof(uint64_t) :
(kpattern_size32 & 0x3) == 0 ?
sizeof(uint32_t) :
(kpattern_size32 & 0x1) == 0 ?
sizeof(uint16_t) : sizeof(uint8_t);
// Program kernels arguments for the fill operation // Program kernels arguments for the fill operation
cl_mem mem = as_cl<amd::Memory>(memory.owner()); cl_mem mem = as_cl<amd::Memory>(memory.owner());
if (alignment == sizeof(uint64_t)) { if (alignment == 2 * sizeof(uint64_t)) {
setArgument(kernels_[fillType], 0, sizeof(cl_mem), nullptr); setArgument(kernels_[kFillType], 0, sizeof(cl_mem), nullptr);
setArgument(kernels_[fillType], 1, sizeof(cl_mem), nullptr); setArgument(kernels_[kFillType], 1, sizeof(cl_mem), nullptr);
setArgument(kernels_[fillType], 2, sizeof(cl_mem), nullptr); setArgument(kernels_[kFillType], 2, sizeof(cl_mem), nullptr);
setArgument(kernels_[fillType], 3, sizeof(cl_mem), &mem); setArgument(kernels_[kFillType], 3, sizeof(cl_mem), nullptr);
setArgument(kernels_[kFillType], 4, sizeof(cl_mem), &mem);
} else if (alignment == sizeof(uint64_t)) {
setArgument(kernels_[kFillType], 0, sizeof(cl_mem), nullptr);
setArgument(kernels_[kFillType], 1, sizeof(cl_mem), nullptr);
setArgument(kernels_[kFillType], 2, sizeof(cl_mem), nullptr);
setArgument(kernels_[kFillType], 3, sizeof(cl_mem), &mem);
setArgument(kernels_[kFillType], 4, sizeof(cl_mem), nullptr);
} else if (alignment == sizeof(uint32_t)) { } else if (alignment == sizeof(uint32_t)) {
setArgument(kernels_[fillType], 0, sizeof(cl_mem), nullptr); setArgument(kernels_[kFillType], 0, sizeof(cl_mem), nullptr);
setArgument(kernels_[fillType], 1, sizeof(cl_mem), nullptr); setArgument(kernels_[kFillType], 1, sizeof(cl_mem), nullptr);
setArgument(kernels_[fillType], 2, sizeof(cl_mem), &mem); setArgument(kernels_[kFillType], 2, sizeof(cl_mem), &mem);
setArgument(kernels_[fillType], 3, sizeof(cl_mem), nullptr); setArgument(kernels_[kFillType], 3, sizeof(cl_mem), nullptr);
setArgument(kernels_[kFillType], 4, sizeof(cl_mem), nullptr);
} else if (alignment == sizeof(uint16_t)) { } else if (alignment == sizeof(uint16_t)) {
setArgument(kernels_[fillType], 0, sizeof(cl_mem), nullptr); setArgument(kernels_[kFillType], 0, sizeof(cl_mem), nullptr);
setArgument(kernels_[fillType], 1, sizeof(cl_mem), &mem); setArgument(kernels_[kFillType], 1, sizeof(cl_mem), &mem);
setArgument(kernels_[fillType], 2, sizeof(cl_mem), nullptr); setArgument(kernels_[kFillType], 2, sizeof(cl_mem), nullptr);
setArgument(kernels_[fillType], 3, sizeof(cl_mem), nullptr); setArgument(kernels_[kFillType], 3, sizeof(cl_mem), nullptr);
setArgument(kernels_[kFillType], 4, sizeof(cl_mem), nullptr);
} else { } else {
setArgument(kernels_[fillType], 0, sizeof(cl_mem), &mem); setArgument(kernels_[kFillType], 0, sizeof(cl_mem), &mem);
setArgument(kernels_[fillType], 1, sizeof(cl_mem), nullptr); setArgument(kernels_[kFillType], 1, sizeof(cl_mem), nullptr);
setArgument(kernels_[fillType], 2, sizeof(cl_mem), nullptr); setArgument(kernels_[kFillType], 2, sizeof(cl_mem), nullptr);
setArgument(kernels_[fillType], 3, sizeof(cl_mem), nullptr); setArgument(kernels_[kFillType], 3, sizeof(cl_mem), nullptr);
setArgument(kernels_[kFillType], 4, sizeof(cl_mem), nullptr);
} }
const size_t localWorkSize = 256;
size_t globalWorkSize =
std::min(dev().settings().limit_blit_wg_ * localWorkSize, kfill_size);
globalWorkSize = amd::alignUp(globalWorkSize, localWorkSize);
auto constBuf = gpu().allocKernArg(kCBSize, kCBAlignment); auto constBuf = gpu().allocKernArg(kCBSize, kCBAlignment);
// If pattern has been expanded, use the expanded pattern, otherwise use the default pattern. // If pattern has been expanded, use the expanded pattern, otherwise use the default pattern.
if (packed_obj.pattern_expanded_) { if (packed_obj.pattern_expanded_) {
memcpy(constBuf, &packed_obj.expanded_pattern_, kpattern_size32); memcpy(constBuf, &packed_obj.expanded_pattern_, kpattern_size);
} else { } else {
memcpy(constBuf, pattern, kpattern_size32); memcpy(constBuf, pattern, kpattern_size);
} }
constexpr bool kDirectVa = true; constexpr bool kDirectVa = true;
setArgument(kernels_[fillType], 4, sizeof(cl_mem), constBuf, 0, nullptr, kDirectVa); setArgument(kernels_[kFillType], 5, sizeof(cl_mem), constBuf, 0, nullptr, kDirectVa);
// Adjust the pattern size in the copy type size
kpattern_size /= alignment;
setArgument(kernels_[kFillType], 6, sizeof(uint32_t), &kpattern_size);
koffset /= alignment; koffset /= alignment;
kpattern_size32 /= alignment; setArgument(kernels_[kFillType], 7, sizeof(koffset), &koffset);
// Calculate max id
setArgument(kernels_[fillType], 5, sizeof(uint32_t), &kpattern_size32); kfill_size = memory.virtualAddress() + (koffset + kfill_size * kpattern_size) * alignment;
setArgument(kernels_[fillType], 6, sizeof(koffset), &koffset); setArgument(kernels_[kFillType], 8, sizeof(kfill_size), &kfill_size);
setArgument(kernels_[fillType], 7, sizeof(kfill_size), &kfill_size); uint32_t next_chunk = globalWorkSize * kpattern_size;
setArgument(kernels_[kFillType], 9, sizeof(uint32_t), &next_chunk);
// Create ND range object for the kernel's execution // Create ND range object for the kernel's execution
amd::NDRangeContainer ndrange(1, globalWorkOffset, &globalWorkSize, &localWorkSize); amd::NDRangeContainer ndrange(1, globalWorkOffset, &globalWorkSize, &localWorkSize);
// Execute the blit // Execute the blit
address parameters = captureArguments(kernels_[fillType]); address parameters = captureArguments(kernels_[kFillType]);
result = gpu().submitKernelInternal(ndrange, *kernels_[fillType], parameters, nullptr); result = gpu().submitKernelInternal(ndrange, *kernels_[kFillType], parameters, nullptr);
releaseArguments(parameters); releaseArguments(parameters);
} }
} }
projects/clr/rocclr/device/rocm/rocdevice.cpp
+4 -1
View File
@@ -769,7 +769,10 @@ bool Device::create() {
pciDeviceId_); pciDeviceId_);
return false; return false;
} }
hsaSettings->limit_blit_wg_ = info().maxComputeUnits_;
if (!flagIsDefault(DEBUG_CLR_LIMIT_BLIT_WG)) {
hsaSettings->limit_blit_wg_ = std::max(DEBUG_CLR_LIMIT_BLIT_WG, 0x1U);
}
amd::Context::Info info = {0}; amd::Context::Info info = {0};
std::vector<amd::Device*> devices; std::vector<amd::Device*> devices;
devices.push_back(this); devices.push_back(this);
projects/clr/rocclr/device/rocm/rocsettings.cpp
+1
View File
@@ -96,6 +96,7 @@ Settings::Settings() {
host_hdp_flush_ = true; host_hdp_flush_ = true;
gwsInitSupported_ = true; gwsInitSupported_ = true;
limit_blit_wg_ = 16;
} }
// ================================================================================================ // ================================================================================================
projects/clr/rocclr/device/rocm/rocsettings.hpp
+1
View File
@@ -76,6 +76,7 @@ class Settings : public device::Settings {
size_t sdmaCopyThreshold_; //!< Use SDMA to copy above this size size_t sdmaCopyThreshold_; //!< Use SDMA to copy above this size
uint32_t hmmFlags_; //!< HMM functionality control flags uint32_t hmmFlags_; //!< HMM functionality control flags
uint32_t limit_blit_wg_; //!< The number of workgroups for blit execution
//! Default constructor //! Default constructor
Settings(); Settings();
projects/clr/rocclr/device/rocm/rocvirtual.cpp
+2 -2
View File
@@ -2828,8 +2828,8 @@ bool VirtualGPU::createVirtualQueue(uint deviceQueueSize)
// Add mask array for AmdAqlWrap slots // Add mask array for AmdAqlWrap slots
allocSize += amd::alignUp(numSlots, DeviceQueueMaskSize) / 8; allocSize += amd::alignUp(numSlots, DeviceQueueMaskSize) / 8;
// Make sure the allocation size aligns with DWORD. // Align size to 64 bytes for more efficient fill operation
allocSize = amd::alignUp(allocSize, sizeof(uint64_t)); allocSize = amd::alignUp(allocSize, 8 * sizeof(uint64_t));
// CL_MEM_ALLOC_HOST_PTR/CL_MEM_READ_WRITE // CL_MEM_ALLOC_HOST_PTR/CL_MEM_READ_WRITE
virtualQueue_ = new (dev().context()) amd::Buffer(dev().context(), CL_MEM_READ_WRITE, allocSize); virtualQueue_ = new (dev().context()) amd::Buffer(dev().context(), CL_MEM_READ_WRITE, allocSize);
projects/clr/rocclr/utils/flags.hpp
+2
View File
@@ -221,6 +221,8 @@ release(uint, ROC_AQL_QUEUE_SIZE, 16384, \
"AQL queue size in AQL packets") \ "AQL queue size in AQL packets") \
release(uint, ROC_SIGNAL_POOL_SIZE, 32, \ release(uint, ROC_SIGNAL_POOL_SIZE, 32, \
"Initial size of HSA signal pool") \ "Initial size of HSA signal pool") \
release(uint, DEBUG_CLR_LIMIT_BLIT_WG, 16, \
"Limit the number of workgroups in blit operations") \
release(bool, ROC_SKIP_KERNEL_ARG_COPY, false, \ release(bool, ROC_SKIP_KERNEL_ARG_COPY, false, \
"If true, then runtime can skip kernel arg copy") \ "If true, then runtime can skip kernel arg copy") \
release(bool, GPU_STREAMOPS_CP_WAIT, false, \ release(bool, GPU_STREAMOPS_CP_WAIT, false, \