[libomptarget][amdgpu][nfc] Update comments

[JonChesterfield]

Summary

[libomptarget][amdgpu][nfc] Update comments

[jdoerfert]

LG

openmp/libomptarget/plugins/amdgpu/impl/data.cpp

 hsa_amd_memory_pool_t get_memory_pool_by_mem_place(atmi_mem_place_t place) {
   ATLProcessor &proc = get_processor_by_mem_place(place);
   return get_memory_pool(proc, place.mem_id);
 }
 
 void register_allocation(void *ptr, size_t size, atmi_mem_place_t place) {
   if (place.dev_type == ATMI_DEVTYPE_CPU)
     allow_access_to_all_gpu_agents(ptr);
-  // TODO(ashwinma): what if one GPU wants to access another GPU?
 }
 
 atmi_status_t Runtime::Malloc(void **ptr, size_t size, atmi_mem_place_t place) {
   atmi_status_t ret = ATMI_STATUS_SUCCESS;
   hsa_amd_memory_pool_t pool = get_memory_pool_by_mem_place(place);
   hsa_status_t err = hsa_amd_memory_pool_allocate(pool, size, 0, ptr);
   ErrorCheck(atmi_malloc, err);
   DEBUG_PRINT("Malloced [%s %d] %p\n",

openmp/libomptarget/plugins/amdgpu/impl/machine.h

 public:
   explicit ATLProcessor(hsa_agent_t agent,
                         atmi_devtype_t type = ATMI_DEVTYPE_ALL)
       : agent_(agent), type_(type) {
     memories_.clear();
   }
   void addMemory(const ATLMemory &p);
   hsa_agent_t agent() const { return agent_; }
-  // TODO(ashwinma): Do we need this or are we building the machine structure
-  // just once in the program?
-  // void removeMemory(ATLMemory &p);
   const std::vector<ATLMemory> &memories() const;
   atmi_devtype_t type() const { return type_; }
 
 protected:
   hsa_agent_t agent_;
   atmi_devtype_t type_;
   std::vector<ATLMemory> memories_;
 };

 hsa_amd_memory_pool_t get_memory_pool(const ATLProcessor &proc,
                                       const int mem_id);
 
 extern ATLMachine g_atl_machine;
 template <typename T> T &get_processor(atmi_place_t place) {
   int dev_id = place.device_id;
   if (dev_id == -1) {
     // user is asking runtime to pick a device
-    // TODO(ashwinma): best device of this type? pick 0 for now
+    // best device of this type? pick 0 for now
     dev_id = 0;
   }
   return g_atl_machine.processors<T>()[dev_id];
 }
 
 #endif // SRC_RUNTIME_INCLUDE_MACHINE_H_

openmp/libomptarget/plugins/amdgpu/impl/rt.h

Environment()
       : max_queue_size_(DEFAULT_MAX_QUEUE_SIZE),
         debug_mode_(DEFAULT_DEBUG_MODE) {
     GetEnvAll();
   }
 
   void GetEnvAll();
 
   int getMaxQueueSize() const { return max_queue_size_; }
-
-  // TODO(ashwinma): int may change to enum if we have more debug modes
   int getDebugMode() const { return debug_mode_; }
-  // TODO(ashwinma): int may change to enum if we have more profile modes
 
 private:
   std::string GetEnv(const char *name) {
     char *env = getenv(name);
     std::string ret;
     if (env) {
       ret = env;
     }
static atmi_status_t RegisterModuleFromMemory(
                                             void *cb_state),
       void *cb_state);
 
   // data
   static atmi_status_t Memcpy(hsa_signal_t, void *, const void *, size_t);
   static atmi_status_t Memfree(void *);
   static atmi_status_t Malloc(void **, size_t, atmi_mem_place_t);
 
-  // environment variables
   int getMaxQueueSize() const { return env_.getMaxQueueSize(); }
-
-  // TODO(ashwinma): int may change to enum if we have more debug modes
   int getDebugMode() const { return env_.getDebugMode(); }
 
 protected:
   Runtime() = default;
   ~Runtime() = default;
   Runtime(const Runtime &) = delete;
   Runtime &operator=(const Runtime &) = delete;
 
 protected:
   // variable to track environment variables
   Environment env_;
 };
 
 } // namespace core
 
 #endif // SRC_RUNTIME_INCLUDE_RT_H_

openmp/libomptarget/plugins/amdgpu/impl/system.cpp

static const std::map<std::string, KernelArgMD::ValueKind> ArgValueKind = {
     {"hidden_default_queue", KernelArgMD::ValueKind::HiddenDefaultQueue},
     {"hidden_completion_action",
      KernelArgMD::ValueKind::HiddenCompletionAction},
     {"hidden_multigrid_sync_arg",
      KernelArgMD::ValueKind::HiddenMultiGridSyncArg},
     {"hidden_hostcall_buffer", KernelArgMD::ValueKind::HiddenHostcallBuffer},
 };
 
-// public variables -- TODO(ashwinma) move these to a runtime object?
+// global variables. TODO: Get rid of these
 atmi_machine_t g_atmi_machine;
 ATLMachine g_atl_machine;
 
 hsa_region_t atl_gpu_kernarg_region;
 std::vector<hsa_amd_memory_pool_t> atl_gpu_kernarg_pools;
 hsa_region_t atl_cpu_kernarg_region;
 
 static std::vector<hsa_executable_t> g_executables;
if (devtype == ATMI_DEVTYPE_ALL || devtype & ATMI_DEVTYPE_GPU) {
     ATMIErrorCheck(GPU context init, atl_init_gpu_context());
   }
 
   atl_set_atmi_initialized();
   return ATMI_STATUS_SUCCESS;
 }
 
 atmi_status_t Runtime::Finalize() {
-  // TODO(ashwinma): Finalize all processors, queues, signals, kernarg memory
-  // regions
   hsa_status_t err;
 
   for (uint32_t i = 0; i < g_executables.size(); i++) {
     err = hsa_executable_destroy(g_executables[i]);
     ErrorCheck(Destroying executable, err);
   }
 
   for (uint32_t i = 0; i < SymbolInfoTable.size(); i++) {
if (kernel_segment_size > 0) {
         msgpack_errors += array_lookup_element(args_array, i, &args_element);
         msgpackErrorCheck(iterate args map in kernel args metadata,
                           msgpack_errors);
 
         msgpack_errors += populate_kernelArgMD(args_element, &lcArg);
         msgpackErrorCheck(iterate args map in kernel args metadata,
                           msgpack_errors);
 
-        // TODO(ashwinma): should the below population actions be done only for
-        // non-implicit args?
         // populate info with sizes and offsets
         info.arg_sizes.push_back(lcArg.size_);
         // v3 has offset field and not align field
         size_t new_offset = lcArg.offset_;
         size_t padding = new_offset - offset;
         offset = new_offset;
         info.arg_offsets.push_back(lcArg.offset_);
         DEBUG_PRINT("Arg[%lu] \"%s\" (%u, %u)\n", i, lcArg.name_.c_str(),

openmp/libomptarget/plugins/amdgpu/src/rtl.cpp

if (state_ptr_size == sizeof(void *)) {
                                                  sizeof(void *), device_id);
       if (err != ATMI_STATUS_SUCCESS) {
         fprintf(stderr, "memcpy install of state_ptr failed\n");
         return NULL;
       }
     }
   }
 
-  // TODO: Check with Guansong to understand the below comment more thoroughly.
   // Here, we take advantage of the data that is appended after img_end to get
   // the symbols' name we need to load. This data consist of the host entries
   // begin and end as well as the target name (see the offloading linker script
   // creation in clang compiler).
 
   // Find the symbols in the module by name. The name can be obtain by
   // concatenating the host entry name with the target name