diff --git a/rocclr/device/blitcl.cpp b/rocclr/device/blitcl.cpp index 6c899c49c5..921f4570bb 100644 --- a/rocclr/device/blitcl.cpp +++ b/rocclr/device/blitcl.cpp @@ -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); diff --git a/rocclr/device/pal/palblit.cpp b/rocclr/device/pal/palblit.cpp index 14b77c83eb..cc4c045397 100644 --- a/rocclr/device/pal/palblit.cpp +++ b/rocclr/device/pal/palblit.cpp @@ -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), ©Size); + setArgument(kernels_[kBlitType], 2, sizeof(copySize), ©Size); - 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); diff --git a/rocclr/device/rocm/rocblit.cpp b/rocclr/device/rocm/rocblit.cpp index a8f2f17d04..ddb458a050 100644 --- a/rocclr/device/rocm/rocblit.cpp +++ b/rocclr/device/rocm/rocblit.cpp @@ -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(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(srcMemory.owner()); - setArgument(kernels_[kBlitType], 0, sizeof(cl_mem), &mem, 0, &srcMemory); - mem = as_cl(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(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), ©Size); + setArgument(kernels_[kBlitType], 2, sizeof(copySize), ©Size); - 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);