AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

SWDEV-455254 - Reduce blit kernels signature

Просмотр исходного кода 
Remove offset from blit kernels, since it can be applied in setup.

Change-Id: I06b585068d68a0ee8e125ddf46a36fccb372f30d
Этот коммит содержится в:
German Andryeyev
2024-04-09 14:27:04 -04:00
родитель 35c80dd482
Коммит 7de7da4016
3 изменённых файлов: 47 добавлений и 114 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
rocclr/device/blitcl.cpp
+14 -13
Просмотреть файл
@@ -48,12 +48,12 @@ const char* BlitLinearSourceCode = BLIT_KERNELS(
extern void __ockl_gws_init(uint nwm1, uint rid);
__kernel void __amd_rocclr_fillBufferAligned(
__global uchar* bufUChar, __global ushort* bufUShort, __global uint* bufUInt,
__global ulong* bufULong, __global ulong2* bufULong2, __constant uchar* pattern,
uint pattern_size, ulong offset, ulong end_ptr, uint next_chunk) {
__global void* buf, __constant uchar* pattern,
uint pattern_size, uint alignment, ulong end_ptr, uint next_chunk) {
int id = get_global_id(0);
long cur_id = offset + id * pattern_size;
if (bufULong2) {
long cur_id = id * pattern_size;
if (alignment == sizeof(ulong2)) {
__global ulong2* bufULong2 = (__global ulong2*)buf;
__global ulong2* element = &bufULong2[cur_id];
__constant ulong2* pt = (__constant ulong2*)pattern;
while ((ulong)element < end_ptr) {
@@ -62,7 +62,8 @@ const char* BlitLinearSourceCode = BLIT_KERNELS(
}
element += next_chunk;
}
} else if (bufULong) {
} else if (alignment == sizeof(ulong)) {
__global ulong* bufULong = (__global ulong*)buf;
__global ulong* element = &bufULong[cur_id];
__constant ulong* pt = (__constant ulong*)pattern;
while ((ulong)element < end_ptr) {
@@ -71,7 +72,8 @@ const char* BlitLinearSourceCode = BLIT_KERNELS(
}
element += next_chunk;
}
} else if (bufUInt) {
} else if (alignment == sizeof(uint)) {
__global uint* bufUInt = (__global uint*)buf;
__global uint* element = &bufUInt[cur_id];
__constant uint* pt = (__constant uint*)pattern;
while ((ulong)element < end_ptr) {
@@ -80,7 +82,8 @@ const char* BlitLinearSourceCode = BLIT_KERNELS(
}
element += next_chunk;
}
} else if (bufUShort) {
} else if (alignment == sizeof(ushort)) {
__global ushort* bufUShort = (__global ushort*)buf;
__global ushort* element = &bufUShort[cur_id];
__constant ushort* pt = (__constant ushort*)pattern;
while ((ulong)element < end_ptr) {
@@ -90,6 +93,7 @@ const char* BlitLinearSourceCode = BLIT_KERNELS(
element += next_chunk;
}
} else {
__global uchar* bufUChar = (__global uchar*)buf;
__global uchar* element = &bufUChar[cur_id];
while ((ulong)element < end_ptr) {
for (uint i = 0; i < pattern_size; ++i) {
@@ -115,15 +119,12 @@ const char* BlitLinearSourceCode = BLIT_KERNELS(
pitch);
}
__kernel void __amd_rocclr_copyBuffer(__global uchar* srcI, __global uchar* dstI,
ulong srcOrigin, ulong dstOrigin, ulong size, uint remainder,
__kernel void __amd_rocclr_copyBuffer(__global uchar* src, __global uchar* dst,
ulong size, uint remainder,
uint aligned_size, ulong end_ptr, uint next_chunk) {
ulong id = get_global_id(0);
ulong id_remainder = id;
__global uchar* src = srcI + srcOrigin;
__global uchar* dst = dstI + dstOrigin;
if (aligned_size == sizeof(ulong2)) {
__global ulong2* srcD = (__global ulong2*)(src);
__global ulong2* dstD = (__global ulong2*)(dst);
rocclr/device/pal/palblit.cpp
+17 -51
Просмотреть файл
@@ -2194,37 +2194,7 @@ bool KernelBlitManager::fillBuffer(device::Memory& memory, const void* pattern,
// Program kernels arguments for the fill operation
Memory* mem = &gpuMem(memory);
if (alignment == 2 * 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), 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);
}
setArgument(kernels_[kFillType], 0, sizeof(cl_mem), &mem, koffset);
const size_t localWorkSize = 256;
size_t globalWorkSize =
std::min(dev().settings().limit_blit_wg_ * localWorkSize, kfill_size);
@@ -2240,20 +2210,20 @@ bool KernelBlitManager::fillBuffer(device::Memory& memory, const void* pattern,
}
gpuCB.unmap(&gpu());
Memory* pGpuCB = &gpuCB;
setArgument(kernels_[kFillType], 5, sizeof(cl_mem), &pGpuCB);
setArgument(kernels_[kFillType], 1, 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);
setArgument(kernels_[kFillType], 2, sizeof(uint32_t), &kpattern_size);
setArgument(kernels_[kFillType], 3, sizeof(alignment), &alignment);
// Calculate max id
uint64_t end_ptr = memory.virtualAddress() +
(koffset + kfill_size * kpattern_size) * alignment;
setArgument(kernels_[kFillType], 8, sizeof(end_ptr), &end_ptr);
uint64_t end_ptr = memory.virtualAddress() + koffset +
kfill_size * kpattern_size * alignment;
setArgument(kernels_[kFillType], 4, sizeof(end_ptr), &end_ptr);
uint32_t next_chunk = globalWorkSize * kpattern_size;
setArgument(kernels_[kFillType], 9, sizeof(uint32_t), &next_chunk);
setArgument(kernels_[kFillType], 5, sizeof(uint32_t), &next_chunk);
// Create ND range object for the kernel's execution
amd::NDRangeContainer ndrange(1, globalWorkOffset, &globalWorkSize, &localWorkSize);
@@ -2297,30 +2267,26 @@ bool KernelBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& ds
// Program kernels arguments for the blit operation
Memory* mem = &gpuMem(srcMemory);
setArgument(kernels_[kBlitType], 0, sizeof(cl_mem), &mem);
mem = &gpuMem(dstMemory);
setArgument(kernels_[kBlitType], 1, sizeof(cl_mem), &mem);
// Program source origin
uint64_t srcOffset = srcOrigin[0];
setArgument(kernels_[kBlitType], 2, sizeof(srcOffset), &srcOffset);
setArgument(kernels_[kBlitType], 0, sizeof(cl_mem), &mem, srcOffset);
mem = &gpuMem(dstMemory);
// Program destinaiton origin
uint64_t dstOffset = dstOrigin[0];
setArgument(kernels_[kBlitType], 3, sizeof(dstOffset), &dstOffset);
setArgument(kernels_[kBlitType], 1, sizeof(cl_mem), &mem, dstOffset);
uint64_t copySize = sizeIn[0];
setArgument(kernels_[kBlitType], 4, sizeof(copySize), &copySize);
setArgument(kernels_[kBlitType], 2, sizeof(copySize), &copySize);
setArgument(kernels_[kBlitType], 5, sizeof(remainder), &remainder);
setArgument(kernels_[kBlitType], 6, sizeof(aligned_size), &aligned_size);
setArgument(kernels_[kBlitType], 3, sizeof(remainder), &remainder);
setArgument(kernels_[kBlitType], 4, sizeof(aligned_size), &aligned_size);
// End pointer is the aligned copy size and destination offset
uint64_t end_ptr = dstMemory.virtualAddress() + dstOffset + sizeIn[0] - remainder;
setArgument(kernels_[kBlitType], 7, sizeof(end_ptr), &end_ptr);
setArgument(kernels_[kBlitType], 5, sizeof(end_ptr), &end_ptr);
uint32_t next_chunk = globalWorkSize;
setArgument(kernels_[kBlitType], 8, sizeof(next_chunk), &next_chunk);
setArgument(kernels_[kBlitType], 6, sizeof(next_chunk), &next_chunk);
// Create ND range object for the kernel's execution
amd::NDRangeContainer ndrange(1, nullptr, &globalWorkSize, &localWorkSize);
rocclr/device/rocm/rocblit.cpp
+16 -50
Просмотреть файл
@@ -2108,37 +2108,7 @@ bool KernelBlitManager::fillBuffer1D(device::Memory& memory, const void* pattern
(kpattern_size & 0x1) == 0 ? sizeof(uint16_t) : sizeof(uint8_t);
// Program kernels arguments for the fill operation
cl_mem mem = as_cl<amd::Memory>(memory.owner());
if (alignment == 2 * 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), 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);
}
setArgument(kernels_[kFillType], 0, sizeof(cl_mem), &mem, koffset);
const size_t localWorkSize = 256;
size_t globalWorkSize =
std::min(dev().settings().limit_blit_wg_ * localWorkSize, kfill_size);
@@ -2153,18 +2123,18 @@ bool KernelBlitManager::fillBuffer1D(device::Memory& memory, const void* pattern
memcpy(constBuf, pattern, kpattern_size);
}
constexpr bool kDirectVa = true;
setArgument(kernels_[kFillType], 5, sizeof(cl_mem), constBuf, 0, nullptr, kDirectVa);
setArgument(kernels_[kFillType], 1, 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;
setArgument(kernels_[kFillType], 7, sizeof(koffset), &koffset);
setArgument(kernels_[kFillType], 2, sizeof(uint32_t), &kpattern_size);
setArgument(kernels_[kFillType], 3, sizeof(alignment), &alignment);
// Calculate max id
kfill_size = memory.virtualAddress() + (koffset + kfill_size * kpattern_size) * alignment;
setArgument(kernels_[kFillType], 8, sizeof(kfill_size), &kfill_size);
kfill_size = memory.virtualAddress() + koffset + kfill_size * kpattern_size * alignment;
setArgument(kernels_[kFillType], 4, sizeof(kfill_size), &kfill_size);
uint32_t next_chunk = globalWorkSize * kpattern_size;
setArgument(kernels_[kFillType], 9, sizeof(uint32_t), &next_chunk);
setArgument(kernels_[kFillType], 5, sizeof(uint32_t), &next_chunk);
// Create ND range object for the kernel's execution
amd::NDRangeContainer ndrange(1, globalWorkOffset, &globalWorkSize, &localWorkSize);
@@ -2351,31 +2321,27 @@ bool KernelBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& ds
// Program kernels arguments for the blit operation
cl_mem mem = as_cl<amd::Memory>(srcMemory.owner());
setArgument(kernels_[kBlitType], 0, sizeof(cl_mem), &mem, 0, &srcMemory);
mem = as_cl<amd::Memory>(dstMemory.owner());
setArgument(kernels_[kBlitType], 1, sizeof(cl_mem), &mem, 0, &dstMemory);
// Program source origin
uint64_t srcOffset = srcOrigin[0];
setArgument(kernels_[kBlitType], 2, sizeof(srcOffset), &srcOffset);
setArgument(kernels_[kBlitType], 0, sizeof(cl_mem), &mem, srcOffset, &srcMemory);
mem = as_cl<amd::Memory>(dstMemory.owner());
// Program destinaiton origin
uint64_t dstOffset = dstOrigin[0];
setArgument(kernels_[kBlitType], 3, sizeof(dstOffset), &dstOffset);
setArgument(kernels_[kBlitType], 1, sizeof(cl_mem), &mem, dstOffset, &dstMemory);
uint64_t copySize = sizeIn[0];
setArgument(kernels_[kBlitType], 4, sizeof(copySize), &copySize);
setArgument(kernels_[kBlitType], 2, sizeof(copySize), &copySize);
setArgument(kernels_[kBlitType], 5, sizeof(remainder), &remainder);
setArgument(kernels_[kBlitType], 6, sizeof(aligned_size), &aligned_size);
setArgument(kernels_[kBlitType], 3, sizeof(remainder), &remainder);
setArgument(kernels_[kBlitType], 4, sizeof(aligned_size), &aligned_size);
// End pointer is the aligned copy size and destination offset
uint64_t end_ptr = dstMemory.virtualAddress() + dstOffset + sizeIn[0] - remainder;
setArgument(kernels_[kBlitType], 7, sizeof(end_ptr), &end_ptr);
setArgument(kernels_[kBlitType], 5, sizeof(end_ptr), &end_ptr);
uint32_t next_chunk = globalWorkSize;
setArgument(kernels_[kBlitType], 8, sizeof(next_chunk), &next_chunk);
setArgument(kernels_[kBlitType], 6, sizeof(next_chunk), &next_chunk);
// Create ND range object for the kernel's execution
amd::NDRangeContainer ndrange(1, nullptr, &globalWorkSize, &localWorkSize);