From be6a06384e17de36f88b9e8e0dd9a144ef12adba Mon Sep 17 00:00:00 2001 From: German Andryeyev Date: Thu, 3 Feb 2022 18:28:56 -0500 Subject: [PATCH] SWDEV-307184 - Add support for the new metadata Metadata in Codeobject version 5 is the extension of CO3 and CO4. Add the detection of the new fields and program them in the setup of the kernel arguments. Change-Id: I27e58df77320ad00f4f16d35912668db803826af --- rocclr/device/devhcmessages.cpp | 6 +- rocclr/device/device.hpp | 12 ++- rocclr/device/devkernel.cpp | 37 +++------ rocclr/device/devkernel.hpp | 61 +++++++++++---- rocclr/device/devprogram.cpp | 5 +- rocclr/device/pal/palkernel.cpp | 111 ++++++++++++++++++++------ rocclr/device/pal/palvirtual.cpp | 14 ++-- rocclr/device/pal/palvirtual.hpp | 37 +-------- rocclr/device/rocm/rocvirtual.cpp | 126 +++++++++++++++++++++--------- rocclr/device/rocm/rocvirtual.hpp | 32 +------- 10 files changed, 262 insertions(+), 179 deletions(-) diff --git a/rocclr/device/devhcmessages.cpp b/rocclr/device/devhcmessages.cpp index c8b9c0df20..34319449dc 100644 --- a/rocclr/device/devhcmessages.cpp +++ b/rocclr/device/devhcmessages.cpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2020 - 2021 Advanced Micro Devices, Inc. +/* Copyright (c) 2020 - 2022 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -58,11 +58,11 @@ struct Message { }; static uint64_t getField(uint64_t desc, uint8_t offset, uint8_t width) { - return (desc >> offset) & ((1UL << width) - 1); + return (desc >> offset) & ((1ULL << width) - 1); } static uint64_t setField(uint64_t desc, uint64_t value, uint8_t offset, uint8_t width) { - uint64_t resetMask = ~(((1UL << width) - 1) << offset); + uint64_t resetMask = ~(((1ULL << width) - 1) << offset); return (desc & resetMask) | (value << offset); } diff --git a/rocclr/device/device.hpp b/rocclr/device/device.hpp index e73ef62452..5a8e9d0301 100644 --- a/rocclr/device/device.hpp +++ b/rocclr/device/device.hpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2008 - 2021 Advanced Micro Devices, Inc. +/* Copyright (c) 2008 - 2022 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -220,6 +220,16 @@ static constexpr const char* OclExtensionsString[] = {"cl_khr_fp64 ", static constexpr int AmdVendor = 0x1002; +template +inline void WriteAqlArgAt(unsigned char* dst, //!< The write pointer to the buffer + T src, //!< The source pointer + uint size, //!< The size in bytes to copy + size_t offset //!< The alignment to follow while writing to the buffer +) { + assert(sizeof(T) == size && "Argument's size mismatches ABI!"); + *(reinterpret_cast(dst + offset)) = src; +} + namespace device { class ClBinary; class BlitManager; diff --git a/rocclr/device/devkernel.cpp b/rocclr/device/devkernel.cpp index 60effb96ed..914811fe70 100644 --- a/rocclr/device/devkernel.cpp +++ b/rocclr/device/devkernel.cpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2008 - 2021 Advanced Micro Devices, Inc. +/* Copyright (c) 2008 - 2022 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -392,23 +392,13 @@ static amd_comgr_status_t populateArgsV3(const amd_comgr_metadata_node_t key, return AMD_COMGR_STATUS_ERROR; } lcArg->info_.oclObject_ = itValueKind->second; - switch (lcArg->info_.oclObject_) { - case amd::KernelParameterDescriptor::MemoryObject: - if (itValueKind->first.compare("dynamic_shared_pointer") == 0) { - lcArg->info_.shared_ = true; - } - break; - case amd::KernelParameterDescriptor::HiddenGlobalOffsetX: - case amd::KernelParameterDescriptor::HiddenGlobalOffsetY: - case amd::KernelParameterDescriptor::HiddenGlobalOffsetZ: - case amd::KernelParameterDescriptor::HiddenPrintfBuffer: - case amd::KernelParameterDescriptor::HiddenHostcallBuffer: - case amd::KernelParameterDescriptor::HiddenDefaultQueue: - case amd::KernelParameterDescriptor::HiddenCompletionAction: - case amd::KernelParameterDescriptor::HiddenMultiGridSync: - case amd::KernelParameterDescriptor::HiddenNone: - lcArg->info_.hidden_ = true; - break; + if (lcArg->info_.oclObject_ == amd::KernelParameterDescriptor::MemoryObject) { + if (itValueKind->first.compare("dynamic_shared_pointer") == 0) { + lcArg->info_.shared_ = true; + } + } else if ((lcArg->info_.oclObject_ >= amd::KernelParameterDescriptor::HiddenNone) && + (lcArg->info_.oclObject_ < amd::KernelParameterDescriptor::HiddenLast)) { + lcArg->info_.hidden_ = true; } } break; @@ -1145,14 +1135,9 @@ bool Kernel::GetAttrCodePropMetadata() { } } break; - case 3: - case 4: { - status = amd::Comgr::iterate_map_metadata(kernelMetaNode, populateKernelMetaV3, - static_cast(this)); - } - break; default: - return false; + status = amd::Comgr::iterate_map_metadata(kernelMetaNode, populateKernelMetaV3, + static_cast(this)); } @@ -1294,7 +1279,7 @@ void Kernel::InitParameters(const amd_comgr_metadata_node_t kernelMD) { if (codeObjectVer() == 2) { status = amd::Comgr::iterate_map_metadata(argsNode, populateArgs, data); } - else if ((codeObjectVer() == 3) || (codeObjectVer() == 4)) { + else if (codeObjectVer() >= 3) { status = amd::Comgr::iterate_map_metadata(argsNode, populateArgsV3, data); } } diff --git a/rocclr/device/devkernel.hpp b/rocclr/device/devkernel.hpp index 511e2116b1..f762e71654 100644 --- a/rocclr/device/devkernel.hpp +++ b/rocclr/device/devkernel.hpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2008 - 2021 Advanced Micro Devices, Inc. +/* Copyright (c) 2008 - 2022 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -36,34 +36,48 @@ class NDRange; struct KernelParameterDescriptor { enum { Value = 0, - HiddenNone = 1, - HiddenGlobalOffsetX = 2, - HiddenGlobalOffsetY = 3, - HiddenGlobalOffsetZ = 4, - HiddenPrintfBuffer = 5, - HiddenDefaultQueue = 6, - HiddenCompletionAction = 7, - MemoryObject = 8, - ReferenceObject = 9, - ValueObject = 10, - ImageObject = 11, - SamplerObject = 12, - QueueObject = 13, + MemoryObject = 1, + ReferenceObject = 2, + ValueObject = 3, + ImageObject = 4, + SamplerObject = 5, + QueueObject = 6, + HiddenNone = 7, + HiddenGlobalOffsetX = 8, + HiddenGlobalOffsetY = 9, + HiddenGlobalOffsetZ = 10, + HiddenPrintfBuffer = 11, + HiddenDefaultQueue = 12, + HiddenCompletionAction = 13, HiddenMultiGridSync = 14, HiddenHostcallBuffer = 15, + HiddenBlockCountX = 16, + HiddenBlockCountY = 17, + HiddenBlockCountZ = 18, + HiddenGroupSizeX = 19, + HiddenGroupSizeY = 20, + HiddenGroupSizeZ = 21, + HiddenRemainderX = 22, + HiddenRemainderY = 23, + HiddenRemainderZ = 24, + HiddenGridDims = 25, + HiddenPrivateBase = 26, + HiddenSharedBase = 27, + HiddenQueuePtr = 28, + HiddenLast = 29 }; clk_value_type_t type_; //!< The parameter's type size_t offset_; //!< Its offset in the parameter's stack size_t size_; //!< Its size in bytes union InfoData { struct { - uint32_t oclObject_ : 4; //!< OCL object type + uint32_t oclObject_ : 6; //!< OCL object type uint32_t readOnly_ : 1; //!< OCL object is read only, applied to memory only uint32_t rawPointer_ : 1; //!< Arguments have a raw GPU VA uint32_t defined_ : 1; //!< The argument was defined by the app uint32_t hidden_ : 1; //!< It's a hidden argument uint32_t shared_ : 1; //!< Dynamic shared memory - uint32_t reserved_ : 3; //!< Reserved + uint32_t reserved_ : 1; //!< Reserved uint32_t arrayIndex_ : 20; //!< Index in the objects array or LDS alignment }; uint32_t allValues_; @@ -262,7 +276,20 @@ static const std::map ArgValueKindV3 = { {"hidden_default_queue", amd::KernelParameterDescriptor::HiddenDefaultQueue}, {"hidden_completion_action", amd::KernelParameterDescriptor::HiddenCompletionAction}, {"hidden_multigrid_sync_arg", amd::KernelParameterDescriptor::HiddenMultiGridSync}, - {"hidden_hostcall_buffer", amd::KernelParameterDescriptor::HiddenHostcallBuffer} + {"hidden_hostcall_buffer", amd::KernelParameterDescriptor::HiddenHostcallBuffer}, + {"hidden_block_count_x", amd::KernelParameterDescriptor::HiddenBlockCountX}, + {"hidden_block_count_y", amd::KernelParameterDescriptor::HiddenBlockCountY}, + {"hidden_block_count_z", amd::KernelParameterDescriptor::HiddenBlockCountZ}, + {"hidden_group_size_x", amd::KernelParameterDescriptor::HiddenGroupSizeX}, + {"hidden_group_size_y", amd::KernelParameterDescriptor::HiddenGroupSizeY}, + {"hidden_group_size_z", amd::KernelParameterDescriptor::HiddenGroupSizeZ}, + {"hidden_remainder_x", amd::KernelParameterDescriptor::HiddenRemainderX}, + {"hidden_remainder_y", amd::KernelParameterDescriptor::HiddenRemainderY}, + {"hidden_remainder_z", amd::KernelParameterDescriptor::HiddenRemainderZ}, + {"hidden_grid_dims", amd::KernelParameterDescriptor::HiddenGridDims}, + {"hidden_private_base", amd::KernelParameterDescriptor::HiddenPrivateBase}, + {"hidden_shared_base", amd::KernelParameterDescriptor::HiddenSharedBase}, + {"hidden_queue_ptr", amd::KernelParameterDescriptor::HiddenQueuePtr} }; static const std::map ArgAccQualV3 = { diff --git a/rocclr/device/devprogram.cpp b/rocclr/device/devprogram.cpp index 3a40c03ae5..eb61f7456e 100644 --- a/rocclr/device/devprogram.cpp +++ b/rocclr/device/devprogram.cpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2008 - 2021 Advanced Micro Devices, Inc. +/* Copyright (c) 2008 - 2022 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -2667,6 +2667,9 @@ bool Program::createKernelMetadataMap(void* binary, size_t binSize) { } else if (minor_version == '1') { ClPrint(amd::LOG_INFO, amd::LOG_CODE, "Using Code Object V4."); codeObjectVer_ = 4; + } else if (minor_version == '2') { + ClPrint(amd::LOG_INFO, amd::LOG_CODE, "Using Code Object V5."); + codeObjectVer_ = 5; } else { ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "Unknown code object metadata minor version [%s.%s].", major_version, minor_version); diff --git a/rocclr/device/pal/palkernel.cpp b/rocclr/device/pal/palkernel.cpp index 8fe5b64fcc..e1c748cb38 100644 --- a/rocclr/device/pal/palkernel.cpp +++ b/rocclr/device/pal/palkernel.cpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. +/* Copyright (c) 2015 - 2022 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -269,6 +269,8 @@ hsa_kernel_dispatch_packet_t* HSAILKernel::loadArguments(VirtualGPU& gpu, const const_address params, size_t ldsAddress, uint64_t vmDefQueue, uint64_t* vmParentWrap) const { + // Provide private and local heap addresses + static constexpr uint AddressShift = LP64_SWITCH(0, 32); const_address parameters = params; uint64_t argList; address aqlArgBuf = gpu.managedBuffer().reserve( @@ -291,34 +293,36 @@ hsa_kernel_dispatch_packet_t* HSAILKernel::loadArguments(VirtualGPU& gpu, const // If signatures don't match, then patch the parameters if (signature.version() != kernel.signature().version()) { - WriteAqlArgAt(aqlArgBuf, parameters, signature.paramsSize() - signature.at(0).offset_, - signature.at(0).offset_); + memcpy(aqlArgBuf + signature.at(0).offset_, parameters, + signature.paramsSize() - signature.at(0).offset_); parameters = aqlArgBuf; } + amd::NDRange local(sizes.local()); + const amd::NDRange& global = sizes.global(); + + // Check if runtime has to find local workgroup size + FindLocalWorkSize(sizes.dimensions(), sizes.global(), local); + + address hidden_arguments = const_cast
(parameters); + // Check if runtime has to setup hidden arguments for (uint32_t i = signature.numParameters(); i < signature.numParametersAll(); ++i) { - const auto it = signature.at(i); - size_t offset; + const auto& it = signature.at(i); switch (it.info_.oclObject_) { case amd::KernelParameterDescriptor::HiddenNone: - // void* zero = 0; - // WriteAqlArgAt(const_cast
(parameters), &zero, it.size_, it.offset_); break; case amd::KernelParameterDescriptor::HiddenGlobalOffsetX: - offset = sizes.offset()[0]; - WriteAqlArgAt(const_cast
(parameters), &offset, it.size_, it.offset_); + WriteAqlArgAt(hidden_arguments, sizes.offset()[0], it.size_, it.offset_); break; case amd::KernelParameterDescriptor::HiddenGlobalOffsetY: if (sizes.dimensions() >= 2) { - offset = sizes.offset()[1]; - WriteAqlArgAt(const_cast
(parameters), &offset, it.size_, it.offset_); + WriteAqlArgAt(hidden_arguments, sizes.offset()[1], it.size_, it.offset_); } break; case amd::KernelParameterDescriptor::HiddenGlobalOffsetZ: if (sizes.dimensions() >= 3) { - offset = sizes.offset()[2]; - WriteAqlArgAt(const_cast
(parameters), &offset, it.size_, it.offset_); + WriteAqlArgAt(hidden_arguments, sizes.offset()[2], it.size_, it.offset_); } break; case amd::KernelParameterDescriptor::HiddenPrintfBuffer: @@ -328,38 +332,101 @@ hsa_kernel_dispatch_packet_t* HSAILKernel::loadArguments(VirtualGPU& gpu, const // and set the fourth argument as the printf_buffer pointer size_t bufferPtr = static_cast(gpu.printfDbgHSA().dbgBuffer()->vmAddress()); gpu.addVmMemory(gpu.printfDbgHSA().dbgBuffer()); - WriteAqlArgAt(const_cast
(parameters), &bufferPtr, it.size_, it.offset_); + WriteAqlArgAt(hidden_arguments, bufferPtr, it.size_, it.offset_); } break; case amd::KernelParameterDescriptor::HiddenHostcallBuffer: if (amd::IS_HIP) { - auto buffer = gpu.getOrCreateHostcallBuffer(); + uintptr_t buffer = reinterpret_cast(gpu.getOrCreateHostcallBuffer()); if (!buffer) { ClPrint(amd::LOG_ERROR, amd::LOG_KERN, "Kernel expects a hostcall buffer, but none found"); } assert(it.size_ == sizeof(buffer) && "check the sizes"); - WriteAqlArgAt(const_cast
(parameters), &buffer, it.size_, it.offset_); + WriteAqlArgAt(hidden_arguments, buffer, it.size_, it.offset_); } break; case amd::KernelParameterDescriptor::HiddenDefaultQueue: if (vmDefQueue != 0) { - WriteAqlArgAt(const_cast
(parameters), &vmDefQueue, it.size_, it.offset_); + WriteAqlArgAt(hidden_arguments, vmDefQueue, it.size_, it.offset_); } break; case amd::KernelParameterDescriptor::HiddenCompletionAction: if (*vmParentWrap != 0) { - WriteAqlArgAt(const_cast
(parameters), vmParentWrap, it.size_, it.offset_); + WriteAqlArgAt(hidden_arguments, *vmParentWrap, it.size_, it.offset_); } break; case amd::KernelParameterDescriptor::HiddenMultiGridSync: break; + case amd::KernelParameterDescriptor::HiddenBlockCountX: + WriteAqlArgAt(hidden_arguments, static_cast(global[0] / local[0]), + it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenBlockCountY: + if (sizes.dimensions() >= 2) { + WriteAqlArgAt(hidden_arguments, static_cast(global[1] / local[1]), + it.size_, it.offset_); + } + break; + case amd::KernelParameterDescriptor::HiddenBlockCountZ: + if (sizes.dimensions() >= 3) { + WriteAqlArgAt(hidden_arguments, static_cast(global[2] / local[2]), + it.size_, it.offset_); + } + break; + case amd::KernelParameterDescriptor::HiddenGroupSizeX: + WriteAqlArgAt(hidden_arguments, static_cast(local[0]), it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenGroupSizeY: + if (sizes.dimensions() >= 2) { + WriteAqlArgAt(hidden_arguments, static_cast(local[1]), it.size_, it.offset_); + } + break; + case amd::KernelParameterDescriptor::HiddenGroupSizeZ: + if (sizes.dimensions() >= 3) { + WriteAqlArgAt(hidden_arguments, static_cast(local[2]), it.size_, it.offset_); + } + break; + case amd::KernelParameterDescriptor::HiddenRemainderX: + WriteAqlArgAt(hidden_arguments, static_cast(global[0] % local[0]), + it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenRemainderY: + if (sizes.dimensions() >= 2) { + WriteAqlArgAt(hidden_arguments, static_cast(global[1] % local[1]), + it.size_, it.offset_); + } + break; + case amd::KernelParameterDescriptor::HiddenRemainderZ: + if (sizes.dimensions() >= 3) { + WriteAqlArgAt(hidden_arguments, static_cast(global[2] % local[2]), + it.size_, it.offset_); + } + break; + case amd::KernelParameterDescriptor::HiddenGridDims: + WriteAqlArgAt(hidden_arguments, static_cast(sizes.dimensions()), + it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenPrivateBase: + WriteAqlArgAt(hidden_arguments, + (gpu.dev().properties().gpuMemoryProperties.privateApertureBase >> AddressShift), + it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenSharedBase: + WriteAqlArgAt(hidden_arguments, + (gpu.dev().properties().gpuMemoryProperties.sharedApertureBase >> AddressShift), + it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenQueuePtr: + // @note: It's not a real AQL queue + WriteAqlArgAt(hidden_arguments, gpu.hsaQueueMem()->vmAddress(), it.size_, it.offset_); + break; } } // Load all kernel arguments if (signature.version() == kernel.signature().version()) { - WriteAqlArgAt(aqlArgBuf, parameters, argsBufferSize(), 0); + memcpy(aqlArgBuf, parameters, argsBufferSize()); } // Note: In a case of structs the size won't match, @@ -371,12 +438,6 @@ hsa_kernel_dispatch_packet_t* HSAILKernel::loadArguments(VirtualGPU& gpu, const hsa_kernel_dispatch_packet_t* hsaDisp = reinterpret_cast(gpu.cb(0)->SysMemCopy()); - amd::NDRange local(sizes.local()); - const amd::NDRange& global = sizes.global(); - - // Check if runtime has to find local workgroup size - FindLocalWorkSize(sizes.dimensions(), sizes.global(), local); - constexpr uint16_t kDispatchPacketHeader = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) | (1 << HSA_PACKET_HEADER_BARRIER) | diff --git a/rocclr/device/pal/palvirtual.cpp b/rocclr/device/pal/palvirtual.cpp index 4e4472c2a6..fad4b874d3 100644 --- a/rocclr/device/pal/palvirtual.cpp +++ b/rocclr/device/pal/palvirtual.cpp @@ -2183,7 +2183,7 @@ void VirtualGPU::PrintChildren(const HSAILKernel& hsaKernel, VirtualGPU* gpuDefQ switch (it.info_.oclObject_) { case amd::KernelParameterDescriptor::HiddenNone: // void* zero = 0; - // WriteAqlArgAt(const_cast
(parameters), &zero, it.size_, it.offset_); + // WriteAqlArgAt(const_cast
(parameters), zero, it.size_, it.offset_); break; case amd::KernelParameterDescriptor::HiddenGlobalOffsetX: extraArgName = "Offset0: "; @@ -3427,7 +3427,7 @@ bool VirtualGPU::processMemObjectsHSA(const amd::Kernel& kernel, const_address p // Save the original LDS size uint64_t ldsSize = *reinterpret_cast(params + desc.offset_); // Patch the LDS address in the original arguments with an LDS address(offset) - WriteAqlArgAt(const_cast
(params), &ldsAddress, desc.size_, desc.offset_); + WriteAqlArgAt(const_cast
(params), ldsAddress, desc.size_, desc.offset_); // Add the original size ldsAddress += ldsSize; } else { @@ -3435,7 +3435,7 @@ bool VirtualGPU::processMemObjectsHSA(const amd::Kernel& kernel, const_address p uint32_t ldsSize = *reinterpret_cast(params + desc.offset_); // Patch the LDS address in the original arguments with an LDS address(offset) uint32_t ldsAddr = ldsAddress; - WriteAqlArgAt(const_cast
(params), &ldsAddr, desc.size_, desc.offset_); + WriteAqlArgAt(const_cast
(params), ldsAddr, desc.size_, desc.offset_); // Add the original size ldsAddress += ldsSize; } @@ -3498,7 +3498,7 @@ bool VirtualGPU::processMemObjectsHSA(const amd::Kernel& kernel, const_address p addBarrier(RgpSqqtBarrierReason::MemDependency); // Use backing store SRD as the replacment uint64_t srd = imageBuffer->CopyImageBuffer()->hwSrd(); - WriteAqlArgAt(const_cast
(params), &srd, sizeof(srd), desc.offset_); + WriteAqlArgAt(const_cast
(params), srd, sizeof(srd), desc.offset_); // Add backing store image to the list of memory handles addVmMemory(imageBuffer->CopyImageBuffer()); // If it's not a read only resource, then runtime has to write back @@ -3517,7 +3517,7 @@ bool VirtualGPU::processMemObjectsHSA(const amd::Kernel& kernel, const_address p uint64_t srd = cb(1)->UploadDataToHw(gpuMem->hwState(), HsaImageObjectSize); // Then use a pointer in aqlArgBuffer to CB1 // Patch the GPU VA address in the original arguments - WriteAqlArgAt(const_cast
(params), &srd, sizeof(srd), desc.offset_); + WriteAqlArgAt(const_cast
(params), srd, sizeof(srd), desc.offset_); addVmMemory(cb(1)->ActiveMemory()); } else { srdResource = true; @@ -3537,7 +3537,7 @@ bool VirtualGPU::processMemObjectsHSA(const amd::Kernel& kernel, const_address p // Then use a pointer in aqlArgBuffer to CB1 const auto it = hsaKernel.patch().find(desc.offset_); // Patch the GPU VA address in the original arguments - WriteAqlArgAt(const_cast
(params), &gpuPtr, sizeof(size_t), it->second); + WriteAqlArgAt(const_cast
(params), gpuPtr, sizeof(size_t), it->second); addVmMemory(cb(1)->ActiveMemory()); } } else if (desc.type_ == T_SAMPLER) { @@ -3558,7 +3558,7 @@ bool VirtualGPU::processMemObjectsHSA(const amd::Kernel& kernel, const_address p vmQueue = vQueue()->vmAddress(); } // Patch the GPU VA address in the original arguments - WriteAqlArgAt(const_cast
(params), &vmQueue, sizeof(vmQueue), desc.offset_); + WriteAqlArgAt(const_cast
(params), vmQueue, sizeof(vmQueue), desc.offset_); break; } } diff --git a/rocclr/device/pal/palvirtual.hpp b/rocclr/device/pal/palvirtual.hpp index 64ee66a3e7..1097fdcfc4 100644 --- a/rocclr/device/pal/palvirtual.hpp +++ b/rocclr/device/pal/palvirtual.hpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. +/* Copyright (c) 2015 - 2022 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -696,17 +696,14 @@ class VirtualGPU : public device::VirtualDevice { inline void VirtualGPU::logVmMemory(const std::string name, const Memory* memory) { if (PAL_EMBED_KERNEL_MD || (AMD_LOG_LEVEL >= amd::LOG_INFO)) { char buf[256]; - sprintf(buf, - "%s = ptr:[%p-%p] obj:[%p-%p]", - name.c_str(), + sprintf(buf, "%s = ptr:[%p-%p] size:[%lld] heap[%d]", name.c_str(), reinterpret_cast(memory->vmAddress()), reinterpret_cast(memory->vmAddress() + memory->size()), - reinterpret_cast(memory->iMem()->Desc().gpuVirtAddr), - reinterpret_cast(memory->iMem()->Desc().gpuVirtAddr + memory->iMem()->Desc().size)); + memory->iMem()->Desc().size, memory->iMem()->Desc().heaps[0]); if (PAL_EMBED_KERNEL_MD) { iCmd()->CmdCommentString(buf); } - LogPrintfInfo("%s threadId : %zx\n", buf, std::this_thread::get_id()); + LogPrintfInfo("%s", buf); } } @@ -735,31 +732,5 @@ template uint VirtualGPU::Queue::submit(bool forceFlus return id; } -template -inline void WriteAqlArgAt(unsigned char* dst, //!< The write pointer to the buffer - const T* src, //!< The source pointer - uint size, //!< The size in bytes to copy - size_t offset //!< The alignment to follow while writing to the buffer -) { - memcpy(dst + offset, src, size); -} - -template <> -inline void WriteAqlArgAt(unsigned char* dst, //!< The write pointer to the buffer - const uint32_t* src, //!< The source pointer - uint size, //!< The size in bytes to copy - size_t offset //!< The alignment to follow while writing to the buffer -) { - *(reinterpret_cast(dst + offset)) = *src; -} - -template <> -inline void WriteAqlArgAt(unsigned char* dst, //!< The write pointer to the buffer - const uint64_t* src, //!< The source pointer - uint size, //!< The size in bytes to copy - size_t offset //!< The alignment to follow while writing to the buffer -) { - *(reinterpret_cast(dst + offset)) = *src; -} /*@}*/ // namespace pal } // namespace pal diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp index 919334b186..e4247a091f 100644 --- a/rocclr/device/rocm/rocvirtual.cpp +++ b/rocclr/device/rocm/rocvirtual.cpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2013 - 2021 Advanced Micro Devices, Inc. +/* Copyright (c) 2013 - 2022 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -34,6 +34,7 @@ #include "utils/debug.hpp" #include "os/os.hpp" #include "amd_hsa_kernel_code.h" +#include "amd_hsa_queue.h" #include #include @@ -643,7 +644,7 @@ bool VirtualGPU::processMemObjects(const amd::Kernel& kernel, const_address para // Save the original LDS size uint64_t ldsSize = *reinterpret_cast(params + desc.offset_); // Patch the LDS address in the original arguments with an LDS address(offset) - WriteAqlArgAt(const_cast
(params), &ldsAddress, desc.size_, desc.offset_); + WriteAqlArgAt(const_cast
(params), ldsAddress, desc.size_, desc.offset_); // Add the original size ldsAddress += ldsSize; } else { @@ -651,7 +652,7 @@ bool VirtualGPU::processMemObjects(const amd::Kernel& kernel, const_address para uint32_t ldsSize = *reinterpret_cast(params + desc.offset_); // Patch the LDS address in the original arguments with an LDS address(offset) uint32_t ldsAddr = ldsAddress; - WriteAqlArgAt(const_cast
(params), &ldsAddr, desc.size_, desc.offset_); + WriteAqlArgAt(const_cast
(params), ldsAddr, desc.size_, desc.offset_); // Add the original size ldsAddress += ldsSize; } @@ -700,7 +701,7 @@ bool VirtualGPU::processMemObjects(const amd::Kernel& kernel, const_address para const uint64_t image_srd = image->getHsaImageObject().handle; assert(amd::isMultipleOf(image_srd, sizeof(image_srd))); - WriteAqlArgAt(const_cast
(params), &image_srd, sizeof(image_srd), desc.offset_); + WriteAqlArgAt(const_cast
(params), image_srd, sizeof(image_srd), desc.offset_); // Check if synchronization has to be performed if (image->CopyImageBuffer() != nullptr) { @@ -717,7 +718,7 @@ bool VirtualGPU::processMemObjects(const amd::Kernel& kernel, const_address para setAqlHeader(dispatchPacketHeader_); // Use backing store SRD as the replacment const uint64_t srd = devCpImg->getHsaImageObject().handle; - WriteAqlArgAt(const_cast
(params), &srd, sizeof(srd), desc.offset_); + WriteAqlArgAt(const_cast
(params), srd, sizeof(srd), desc.offset_); // If it's not a read only resource, then runtime has to write back if (!desc.info_.readOnly_) { @@ -738,7 +739,7 @@ bool VirtualGPU::processMemObjects(const amd::Kernel& kernel, const_address para return false; } uint64_t vqVA = getVQVirtualAddress(); - WriteAqlArgAt(const_cast
(params), &vqVA, sizeof(vqVA), desc.offset_); + WriteAqlArgAt(const_cast
(params), vqVA, sizeof(vqVA), desc.offset_); } else if (desc.type_ == T_VOID) { const_address srcArgPtr = params + desc.offset_; @@ -746,7 +747,7 @@ bool VirtualGPU::processMemObjects(const amd::Kernel& kernel, const_address para void* mem = allocKernArg(desc.size_, 128); memcpy(mem, srcArgPtr, desc.size_); const auto it = hsaKernel.patch().find(desc.offset_); - WriteAqlArgAt(const_cast
(params), &mem, sizeof(void*), it->second); + WriteAqlArgAt(const_cast
(params), mem, sizeof(void*), it->second); } ClPrint(amd::LOG_INFO, amd::LOG_KERN, "Arg%d: %s %s = val:%lld", i, desc.typeName_.c_str(), desc.name_.c_str(), @@ -760,7 +761,7 @@ bool VirtualGPU::processMemObjects(const amd::Kernel& kernel, const_address para device::Sampler* devSampler = sampler->getDeviceSampler(dev()); uint64_t sampler_srd = devSampler->hwSrd(); - WriteAqlArgAt(const_cast
(params), &sampler_srd, sizeof(sampler_srd), desc.offset_); + WriteAqlArgAt(const_cast
(params), sampler_srd, sizeof(sampler_srd), desc.offset_); } } @@ -2734,55 +2735,48 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const ClPrint(amd::LOG_INFO, amd::LOG_KERN, "ShaderName : %s", gpuKernel.name().c_str()); + amd::NDRange local(sizes.local()); + address hidden_arguments = const_cast
(parameters); + // Check if runtime has to setup hidden arguments for (uint32_t i = signature.numParameters(); i < signature.numParametersAll(); ++i) { - const auto it = signature.at(i); - size_t offset; + const auto& it = signature.at(i); switch (it.info_.oclObject_) { case amd::KernelParameterDescriptor::HiddenNone: break; case amd::KernelParameterDescriptor::HiddenGlobalOffsetX: { - offset = newOffset[0]; - assert(it.size_ == sizeof(offset) && "check the sizes"); - WriteAqlArgAt(const_cast
(parameters), &offset, it.size_, it.offset_); + WriteAqlArgAt(hidden_arguments, newOffset[0], it.size_, it.offset_); break; } case amd::KernelParameterDescriptor::HiddenGlobalOffsetY: { if (sizes.dimensions() >= 2) { - offset = newOffset[1]; - assert(it.size_ == sizeof(offset) && "check the sizes"); - WriteAqlArgAt(const_cast
(parameters), &offset, it.size_, it.offset_); + WriteAqlArgAt(hidden_arguments, newOffset[1], it.size_, it.offset_); } break; } case amd::KernelParameterDescriptor::HiddenGlobalOffsetZ: { if (sizes.dimensions() >= 3) { - offset = newOffset[2]; - assert(it.size_ == sizeof(offset) && "check the sizes"); - WriteAqlArgAt(const_cast
(parameters), &offset, it.size_, it.offset_); + WriteAqlArgAt(hidden_arguments, newOffset[2], it.size_, it.offset_); } break; } case amd::KernelParameterDescriptor::HiddenPrintfBuffer: { - address bufferPtr = printfDbg()->dbgBuffer(); - if (printfEnabled && - // and printf buffer was allocated - (bufferPtr != nullptr)) { - assert(it.size_ == sizeof(bufferPtr) && "check the sizes"); - WriteAqlArgAt(const_cast
(parameters), &bufferPtr, it.size_, it.offset_); + uintptr_t bufferPtr = reinterpret_cast(printfDbg()->dbgBuffer()); + if (printfEnabled && !bufferPtr) { + WriteAqlArgAt(hidden_arguments, bufferPtr, it.size_, it.offset_); } break; } case amd::KernelParameterDescriptor::HiddenHostcallBuffer: { if (amd::IS_HIP) { - auto buffer = roc_device_.getOrCreateHostcallBuffer(gpu_queue_, coopGroups, cuMask_); + uintptr_t buffer = reinterpret_cast( + roc_device_.getOrCreateHostcallBuffer(gpu_queue_, coopGroups, cuMask_)); if (!buffer) { ClPrint(amd::LOG_ERROR, amd::LOG_KERN, "Kernel expects a hostcall buffer, but none found"); return false; } - assert(it.size_ == sizeof(buffer) && "check the sizes"); - WriteAqlArgAt(const_cast
(parameters), &buffer, it.size_, it.offset_); + WriteAqlArgAt(hidden_arguments, buffer, it.size_, it.offset_); } break; } @@ -2795,20 +2789,21 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const } vqVA = getVQVirtualAddress(); } - WriteAqlArgAt(const_cast
(parameters), &vqVA, it.size_, it.offset_); + WriteAqlArgAt(hidden_arguments, vqVA, it.size_, it.offset_); break; } case amd::KernelParameterDescriptor::HiddenCompletionAction: { uint64_t spVA = 0; if (nullptr != schedulerParam_) { Memory* schedulerMem = dev().getRocMemory(schedulerParam_); - AmdAqlWrap* wrap = reinterpret_cast(reinterpret_cast(schedulerParam_->getHostMem()) + sizeof(SchedulerParam)); + AmdAqlWrap* wrap = reinterpret_cast( + reinterpret_cast(schedulerParam_->getHostMem()) + sizeof(SchedulerParam)); memset(wrap, 0, sizeof(AmdAqlWrap)); wrap->state = AQL_WRAP_DONE; spVA = reinterpret_cast(schedulerMem->getDeviceMemory()) + sizeof(SchedulerParam); } - WriteAqlArgAt(const_cast
(parameters), &spVA, it.size_, it.offset_); + WriteAqlArgAt(hidden_arguments, spVA, it.size_, it.offset_); break; } case amd::KernelParameterDescriptor::HiddenMultiGridSync: { @@ -2829,20 +2824,82 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const // Update GPU address for grid sync info. Use the offset adjustment for the right location gridSync = reinterpret_cast(syncInfo); } - WriteAqlArgAt(const_cast
(parameters), &gridSync, it.size_, it.offset_); + WriteAqlArgAt(hidden_arguments, gridSync, it.size_, it.offset_); break; } + case amd::KernelParameterDescriptor::HiddenBlockCountX: + WriteAqlArgAt(hidden_arguments, static_cast(newGlobalSize[0] / local[0]), + it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenBlockCountY: + if (sizes.dimensions() >= 2) { + WriteAqlArgAt(hidden_arguments, static_cast(newGlobalSize[1] / local[1]), + it.size_, it.offset_); + } + break; + case amd::KernelParameterDescriptor::HiddenBlockCountZ: + if (sizes.dimensions() >= 3) { + WriteAqlArgAt(hidden_arguments, static_cast(newGlobalSize[2] / local[2]), + it.size_, it.offset_); + } + break; + case amd::KernelParameterDescriptor::HiddenGroupSizeX: + WriteAqlArgAt(hidden_arguments, static_cast(local[0]), it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenGroupSizeY: + if (sizes.dimensions() >= 2) { + WriteAqlArgAt(hidden_arguments, static_cast(local[1]), it.size_, it.offset_); + } + break; + case amd::KernelParameterDescriptor::HiddenGroupSizeZ: + if (sizes.dimensions() >= 3) { + WriteAqlArgAt(hidden_arguments, static_cast(local[2]), it.size_, it.offset_); + } + break; + case amd::KernelParameterDescriptor::HiddenRemainderX: + WriteAqlArgAt(hidden_arguments, static_cast(newGlobalSize[0] % local[0]), + it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenRemainderY: + if (sizes.dimensions() >= 2) { + WriteAqlArgAt(hidden_arguments, static_cast(newGlobalSize[1] % local[1]), + it.size_, it.offset_); + } + break; + case amd::KernelParameterDescriptor::HiddenRemainderZ: + if (sizes.dimensions() >= 3) { + WriteAqlArgAt(hidden_arguments, static_cast(newGlobalSize[2] % local[2]), + it.size_, it.offset_); + } + break; + case amd::KernelParameterDescriptor::HiddenGridDims: + WriteAqlArgAt(hidden_arguments, static_cast(sizes.dimensions()), + it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenPrivateBase: + WriteAqlArgAt(hidden_arguments, + reinterpret_cast(gpu_queue_)->private_segment_aperture_base_hi, + it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenSharedBase: + WriteAqlArgAt(hidden_arguments, + reinterpret_cast(gpu_queue_)->group_segment_aperture_base_hi, + it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenQueuePtr: + WriteAqlArgAt(hidden_arguments, gpu_queue_, it.size_, it.offset_); + break; } } - address argBuffer = const_cast
(parameters); + address argBuffer = hidden_arguments; // Find all parameters for the current kernel if (!kernel.parameters().deviceKernelArgs() || gpuKernel.isInternalKernel()) { // Allocate buffer to hold kernel arguments argBuffer = reinterpret_cast
(allocKernArg(gpuKernel.KernargSegmentByteSize(), gpuKernel.KernargSegmentAlignment())); // Load all kernel arguments - WriteAqlArgAt(argBuffer, parameters, gpuKernel.KernargSegmentByteSize(), 0); + memcpy(argBuffer, parameters, gpuKernel.KernargSegmentByteSize()); } // Note: In a case of structs the size won't match, @@ -2871,7 +2928,6 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const dispatchPacket.grid_size_y = sizes.dimensions() > 1 ? newGlobalSize[1] : 1; dispatchPacket.grid_size_z = sizes.dimensions() > 2 ? newGlobalSize[2] : 1; - amd::NDRange local(sizes.local()); devKernel->FindLocalWorkSize(sizes.dimensions(), sizes.global(), local); dispatchPacket.workgroup_size_x = sizes.dimensions() > 0 ? local[0] : 1; dispatchPacket.workgroup_size_y = sizes.dimensions() > 1 ? local[1] : 1; diff --git a/rocclr/device/rocm/rocvirtual.hpp b/rocclr/device/rocm/rocvirtual.hpp index e8ad55c413..b780052ac4 100644 --- a/rocclr/device/rocm/rocvirtual.hpp +++ b/rocclr/device/rocm/rocvirtual.hpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2008 - 2021 Advanced Micro Devices, Inc. +/* Copyright (c) 2008 - 2022 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -519,34 +519,4 @@ class VirtualGPU : public device::VirtualDevice { //!< OCL doesn't distinguish diffrent copy types, //!< but ROC profiler expects D2H or H2D detection }; - -template -inline void WriteAqlArgAt( - unsigned char* dst, //!< The write pointer to the buffer - const T* src, //!< The source pointer - uint size, //!< The size in bytes to copy - size_t offset //!< The alignment to follow while writing to the buffer -) { - memcpy(dst + offset, src, size); -} - -template <> -inline void WriteAqlArgAt( - unsigned char* dst, //!< The write pointer to the buffer - const uint32_t* src, //!< The source pointer - uint size, //!< The size in bytes to copy - size_t offset //!< The alignment to follow while writing to the buffer -) { - *(reinterpret_cast(dst + offset)) = *src; -} - -template <> -inline void WriteAqlArgAt( - unsigned char* dst, //!< The write pointer to the buffer - const uint64_t* src, //!< The source pointer - uint size, //!< The size in bytes to copy - size_t offset //!< The alignment to follow while writing to the buffer -) { - *(reinterpret_cast(dst + offset)) = *src; -} }