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


[ROCm/clr commit: be6a06384e]

projects/clr/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);
 }
 

projects/clr/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 <typename T>
+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<T*>(dst + offset)) = src;
+}
+
 namespace device {
 class ClBinary;
 class BlitManager;

projects/clr/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<void*>(this));
-      }
-      break;
     default:
-      return false;
+      status = amd::Comgr::iterate_map_metadata(kernelMetaNode, populateKernelMetaV3,
+                                                static_cast<void*>(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);
       }
     }

projects/clr/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<std::string, uint32_t> 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<std::string, cl_kernel_arg_access_qualifier> ArgAccQualV3 = {

projects/clr/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);

projects/clr/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<address>(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<address>(parameters), &zero, it.size_, it.offset_);
         break;
       case amd::KernelParameterDescriptor::HiddenGlobalOffsetX:
-        offset = sizes.offset()[0];
-        WriteAqlArgAt(const_cast<address>(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<address>(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<address>(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<size_t>(gpu.printfDbgHSA().dbgBuffer()->vmAddress());
           gpu.addVmMemory(gpu.printfDbgHSA().dbgBuffer());
-          WriteAqlArgAt(const_cast<address>(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<uintptr_t>(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<address>(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<address>(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<address>(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<uint32_t>(global[0] / local[0]),
+                      it.size_, it.offset_);
+        break;
+      case amd::KernelParameterDescriptor::HiddenBlockCountY:
+        if (sizes.dimensions() >= 2) {
+          WriteAqlArgAt(hidden_arguments, static_cast<uint32_t>(global[1] / local[1]),
+                        it.size_, it.offset_);
+        }
+        break;
+      case amd::KernelParameterDescriptor::HiddenBlockCountZ:
+        if (sizes.dimensions() >= 3) {
+          WriteAqlArgAt(hidden_arguments, static_cast<uint32_t>(global[2] / local[2]),
+                        it.size_, it.offset_);
+        }
+        break;
+      case amd::KernelParameterDescriptor::HiddenGroupSizeX:
+        WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(local[0]), it.size_, it.offset_);
+        break;
+      case amd::KernelParameterDescriptor::HiddenGroupSizeY:
+        if (sizes.dimensions() >= 2) {
+          WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(local[1]), it.size_, it.offset_);
+        }
+        break;
+      case amd::KernelParameterDescriptor::HiddenGroupSizeZ:
+        if (sizes.dimensions() >= 3) {
+          WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(local[2]), it.size_, it.offset_);
+        }
+        break;
+      case amd::KernelParameterDescriptor::HiddenRemainderX:
+        WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(global[0] % local[0]), 
+                      it.size_, it.offset_);
+        break;
+      case amd::KernelParameterDescriptor::HiddenRemainderY:
+        if (sizes.dimensions() >= 2) {
+          WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(global[1] % local[1]),
+                        it.size_, it.offset_);
+        }
+        break;
+      case amd::KernelParameterDescriptor::HiddenRemainderZ:
+        if (sizes.dimensions() >= 3) {
+          WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(global[2] % local[2]),
+                        it.size_, it.offset_);
+        }
+        break;
+      case amd::KernelParameterDescriptor::HiddenGridDims:
+        WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(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<hsa_kernel_dispatch_packet_t*>(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) |

projects/clr/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<address>(parameters), &zero, it.size_, it.offset_);
+            // WriteAqlArgAt(const_cast<address>(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<const uint64_t*>(params + desc.offset_);
           // Patch the LDS address in the original arguments with an LDS address(offset)
-          WriteAqlArgAt(const_cast<address>(params), &ldsAddress, desc.size_, desc.offset_);
+          WriteAqlArgAt(const_cast<address>(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<const uint32_t*>(params + desc.offset_);
           // Patch the LDS address in the original arguments with an LDS address(offset)
           uint32_t ldsAddr = ldsAddress;
-          WriteAqlArgAt(const_cast<address>(params), &ldsAddr, desc.size_, desc.offset_);
+          WriteAqlArgAt(const_cast<address>(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<address>(params), &srd, sizeof(srd), desc.offset_);
+                WriteAqlArgAt(const_cast<address>(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<address>(params), &srd, sizeof(srd), desc.offset_);
+              WriteAqlArgAt(const_cast<address>(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<address>(params), &gpuPtr, sizeof(size_t), it->second);
+        WriteAqlArgAt(const_cast<address>(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<address>(params), &vmQueue, sizeof(vmQueue), desc.offset_);
+      WriteAqlArgAt(const_cast<address>(params), vmQueue, sizeof(vmQueue), desc.offset_);
       break;
     }
   }

projects/clr/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<void*>(memory->vmAddress()),
             reinterpret_cast<void*>(memory->vmAddress() + memory->size()),
-            reinterpret_cast<void*>(memory->iMem()->Desc().gpuVirtAddr),
-            reinterpret_cast<void*>(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 <bool avoidBarrierSubmit> uint VirtualGPU::Queue::submit(bool forceFlus
   return id;
 }
 
-template <typename T>
-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<uint32_t*>(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<uint64_t*>(dst + offset)) = *src;
-}
 /*@}*/  // namespace pal
 }  // namespace pal

projects/clr/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 <fstream>
 #include <limits>
@@ -643,7 +644,7 @@ bool VirtualGPU::processMemObjects(const amd::Kernel& kernel, const_address para
           // Save the original LDS size
           uint64_t ldsSize = *reinterpret_cast<const uint64_t*>(params + desc.offset_);
           // Patch the LDS address in the original arguments with an LDS address(offset)
-          WriteAqlArgAt(const_cast<address>(params), &ldsAddress, desc.size_, desc.offset_);
+          WriteAqlArgAt(const_cast<address>(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<const uint32_t*>(params + desc.offset_);
           // Patch the LDS address in the original arguments with an LDS address(offset)
           uint32_t ldsAddr = ldsAddress;
-          WriteAqlArgAt(const_cast<address>(params), &ldsAddr, desc.size_, desc.offset_);
+          WriteAqlArgAt(const_cast<address>(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<address>(params), &image_srd, sizeof(image_srd), desc.offset_);
+            WriteAqlArgAt(const_cast<address>(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<address>(params), &srd, sizeof(srd), desc.offset_);
+              WriteAqlArgAt(const_cast<address>(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<address>(params), &vqVA, sizeof(vqVA), desc.offset_);
+      WriteAqlArgAt(const_cast<address>(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<address>(params), &mem, sizeof(void*), it->second);
+        WriteAqlArgAt(const_cast<address>(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<address>(params), &sampler_srd, sizeof(sampler_srd), desc.offset_);
+      WriteAqlArgAt(const_cast<address>(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<address>(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<address>(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<address>(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<address>(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<address>(parameters), &bufferPtr, it.size_, it.offset_);
+          uintptr_t bufferPtr = reinterpret_cast<uintptr_t>(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<uintptr_t>(
+              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<address>(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<address>(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<AmdAqlWrap*>(reinterpret_cast<uint64_t>(schedulerParam_->getHostMem()) + sizeof(SchedulerParam));
+            AmdAqlWrap* wrap = reinterpret_cast<AmdAqlWrap*>(
+                               reinterpret_cast<uint64_t>(schedulerParam_->getHostMem()) + sizeof(SchedulerParam));
             memset(wrap, 0, sizeof(AmdAqlWrap));
             wrap->state = AQL_WRAP_DONE;
 
             spVA = reinterpret_cast<uint64_t>(schedulerMem->getDeviceMemory()) + sizeof(SchedulerParam);
           }
-          WriteAqlArgAt(const_cast<address>(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<uint64_t>(syncInfo);
           }
-          WriteAqlArgAt(const_cast<address>(parameters), &gridSync, it.size_, it.offset_);
+          WriteAqlArgAt(hidden_arguments, gridSync, it.size_, it.offset_);
           break;
         }
+        case amd::KernelParameterDescriptor::HiddenBlockCountX:
+          WriteAqlArgAt(hidden_arguments, static_cast<uint32_t>(newGlobalSize[0] / local[0]),
+                        it.size_, it.offset_);
+          break;
+        case amd::KernelParameterDescriptor::HiddenBlockCountY:
+          if (sizes.dimensions() >= 2) {
+            WriteAqlArgAt(hidden_arguments, static_cast<uint32_t>(newGlobalSize[1] / local[1]),
+                          it.size_, it.offset_);
+          }
+          break;
+        case amd::KernelParameterDescriptor::HiddenBlockCountZ:
+          if (sizes.dimensions() >= 3) {
+            WriteAqlArgAt(hidden_arguments, static_cast<uint32_t>(newGlobalSize[2] / local[2]),
+                          it.size_, it.offset_);
+          }
+          break;
+        case amd::KernelParameterDescriptor::HiddenGroupSizeX:
+          WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(local[0]), it.size_, it.offset_);
+          break;
+        case amd::KernelParameterDescriptor::HiddenGroupSizeY:
+          if (sizes.dimensions() >= 2) {
+            WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(local[1]), it.size_, it.offset_);
+          }
+          break;
+        case amd::KernelParameterDescriptor::HiddenGroupSizeZ:
+          if (sizes.dimensions() >= 3) {
+            WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(local[2]), it.size_, it.offset_);
+          }
+          break;
+        case amd::KernelParameterDescriptor::HiddenRemainderX:
+          WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(newGlobalSize[0] % local[0]), 
+                        it.size_, it.offset_);
+          break;
+        case amd::KernelParameterDescriptor::HiddenRemainderY:
+          if (sizes.dimensions() >= 2) {
+            WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(newGlobalSize[1] % local[1]),
+                          it.size_, it.offset_);
+          }
+          break;
+        case amd::KernelParameterDescriptor::HiddenRemainderZ:
+          if (sizes.dimensions() >= 3) {
+            WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(newGlobalSize[2] % local[2]),
+                          it.size_, it.offset_);
+          }
+          break;
+        case amd::KernelParameterDescriptor::HiddenGridDims:
+          WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(sizes.dimensions()),
+                        it.size_, it.offset_);
+          break;
+        case amd::KernelParameterDescriptor::HiddenPrivateBase:
+          WriteAqlArgAt(hidden_arguments,
+                        reinterpret_cast<amd_queue_t*>(gpu_queue_)->private_segment_aperture_base_hi,
+                        it.size_, it.offset_);
+          break;
+        case amd::KernelParameterDescriptor::HiddenSharedBase:
+          WriteAqlArgAt(hidden_arguments,
+                        reinterpret_cast<amd_queue_t*>(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<address>(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<address>(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;

projects/clr/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 <typename T>
-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<uint32_t*>(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<uint64_t*>(dst + offset)) = *src;
-}
 }