[AMDGPU][Libomptarget] Remove global KernelNameMap

[pdhaliwal]

Summary

KernelNameMap contains entries like "key.kd" => key which clearly
could be replaced by simple logic of removing suffix from the key.

[JonChesterfield]

Let's delete it instead. It maps 'foo' to 'foo.kd'. we can replace that with a function that appends .kd to the string.

[pdhaliwal]

Removed KernelNameMap.

[JonChesterfield]

I like the direction. Could we hold it for a day or so? I'd like to check through the uses of the kernel name to see if there's a missing edge case, or if we can simplify this a step further.

It looks like the msgpack data always contains the foo and the foo.kd strings, under different keys. I wonder if that's something we can rely on the compiler emitting.

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

864

I think this is the earliest point we can check the foo to foo.kd mapping is in place, so that we can error out if it isn't

1036

Can we use the foo.kd form everywhere instead of truncating it?

[t-tye]

In D102691#2767846, @JonChesterfield wrote:

I like the direction. Could we hold it for a day or so? I'd like to check through the uses of the kernel name to see if there's a missing edge case, or if we can simplify this a step further.

It looks like the msgpack data always contains the foo and the foo.kd strings, under different keys. I wonder if that's something we can rely on the compiler emitting.

The HSA ABI metadata is defined in https://llvm.org/docs/AMDGPUUsage.html#code-object-metadata . The kernel .symbol attribute provides the ELF symbol that should be used to access the kernel descriptor used in the AQL packet to specify the kernel to execute. The intent is not to make assumptions about the symbols, but simply use the metadata.

[JonChesterfield]

Indeed, but HSA is not the only piece here. This is tied to some generic openmp data, which I think uses the IR symbol name as the ID. That seems to be .symbol.

So far, above co v2 which is unimplemented, kernel name is derivable by appending .kd, so the suggestion here was to only use one of the two strings.

[JonChesterfield]

Taken a closer look at this. Cut down example msgpack,

".name" : "__omp_vmul_l8",
".symbol" : "__omp_vmul_l8.kd",

Most things use the 'name' version, i.e. without the trailing .kd suffix. The symbol name passed into the openmp runtime is that one, and that's the form used as a key in KernelInfoTable.

The HSA functions expect the suffixed version, and that's the one SYMBOL_INFO_NAME returns.

So what we presently do is:

• Build a map from foo.kd to foo
• Check, obliquely, that the original metadata had both foo.kd and foo present
• Use that map to strip .kd from the result of a HSA call

The logic to drop a trailing .kd therefore looks like the right change to make. We can check the .name and .symbol fields when reading the metadata and error out if they don't match the expected format, and likewise error out if HSA returns a string that doesn't end .kd, and at that point we can rip out the map.

That's close to the patch as written above, but now I've actually worked through the control flow myself as well.

[JonChesterfield]

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

137

these ^ are probably not the ideal structures. We convert char* into std::string at each call site, and map() is expected to be slower than unordered_map for the ~random access lookup.

vector<unordered_map<const char*, info_t>> with the custom equality / hasher for C strings is a reasonable choice.

1036

^that turns out to be worse, truncating is the way to go. Probably cleaner to mutate the name[] buffer than do extra things with strings.

[pdhaliwal]

Added check for .kd in symbol name.

[JonChesterfield]

Couple of style comments above. I think this works as-is, so we have confirmation that the map of strings can go. Which is great!

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

877

I think the above works but could be clearer as:
(kernelName + ".kd" ) != symbolName

kernelName can drop out of scope after this test (probably move it into a nested brace to make that clear), so we could mutate it instead of making a new string to probably save an allocation. But in general, '(foo + ".kd") == bar' seems simpler than the find calls.

927–928

Lets drop the debug print and comment here now that the map is gone

[pdhaliwal]

Simplified if condition.

[JonChesterfield]

Like it, thanks!

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

 };
 
 // global variables. TODO: Get rid of these
 atmi_machine_t g_atmi_machine;
 ATLMachine g_atl_machine;
 
 std::vector<hsa_amd_memory_pool_t> atl_gpu_kernarg_pools;
 
-std::map<std::string, std::string> KernelNameMap;
 std::vector<std::map<std::string, atl_kernel_info_t>> KernelInfoTable;
 std::vector<std::map<std::string, atl_symbol_info_t>> SymbolInfoTable;
 
 bool g_atmi_initialized = false;
 bool g_atmi_hostcall_required = false;
 
 /*
    atlc is all internal global values.
    The structure atl_context_t is defined in atl_internal.h
    Most references will use the global structure prefix atlc.
 */
for (size_t i = 0; i < kernelsSize; i++) {
     if (msgpack_errors != 0) {
       printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
              "element lookup in kernel metadata");
       return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
     }
 
     msgpack_errors += map_lookup_string(element, ".name", &kernelName);
     msgpack_errors += map_lookup_string(element, ".symbol", &symbolName);
     if (msgpack_errors != 0) {
       printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
              "strings lookup in kernel metadata");
       return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
     }
 
+    // Make sure that kernelName + ".kd" == symbolName
+    if ((kernelName + ".kd") != symbolName) {
+      printf("[%s:%d] Kernel name mismatching symbol: %s != %s + .kd\n",
+             __FILE__, __LINE__, symbolName.c_str(), kernelName.c_str());
+      return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
+    }
+
     atl_kernel_info_t info = {0, 0, 0, 0, 0, 0, 0, 0, 0, {}, {}, {}};
 
     uint64_t sgpr_count, vgpr_count, sgpr_spill_count, vgpr_spill_count;
     msgpack_errors += map_lookup_uint64_t(element, ".sgpr_count", &sgpr_count);
     if (msgpack_errors != 0) {
       printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
              "sgpr count metadata lookup in kernel metadata");
       return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
     }
for (size_t i = 0; i < kernelsSize; i++) {
     uint64_t kernel_segment_size;
     msgpack_errors += map_lookup_uint64_t(element, ".kernarg_segment_size",
                                           &kernel_segment_size);
     if (msgpack_errors != 0) {
       printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
              "kernarg segment size metadata lookup in kernel metadata");
       return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
     }
 
-    // create a map from symbol to name
-    DEBUG_PRINT("Kernel symbol %s; Name: %s; Size: %lu\n", symbolName.c_str(),
-                kernelName.c_str(), kernel_segment_size);
-    KernelNameMap[symbolName] = kernelName;
-
     bool hasHiddenArgs = false;
     if (kernel_segment_size > 0) {
       uint64_t argsSize;
       size_t offset = 0;
 
       msgpack::byte_range args_array;
       msgpack_errors +=
           map_lookup_array(element, ".args", &args_array, &argsSize);
       if (msgpack_errors != 0) {
if (type == HSA_SYMBOL_KIND_KERNEL) {
     char *name = reinterpret_cast<char *>(malloc(name_length + 1));
     err = hsa_executable_symbol_get_info(symbol,
                                          HSA_EXECUTABLE_SYMBOL_INFO_NAME, name);
     if (err != HSA_STATUS_SUCCESS) {
       printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
              "Symbol info extraction", get_error_string(err));
       return err;
     }
-    name[name_length] = 0;
+    // remove the suffix .kd from symbol name.
+    name[name_length - 3] = 0;
 
-    if (KernelNameMap.find(std::string(name)) == KernelNameMap.end()) {
-      // did not find kernel name in the kernel map; this can happen only
-      // if the ROCr API for getting symbol info (name) is different from
-      // the comgr method of getting symbol info
-      return HSA_STATUS_ERROR;
-    }
     atl_kernel_info_t info;
-    std::string kernelName = KernelNameMap[std::string(name)];
+    std::string kernelName(name);
     // by now, the kernel info table should already have an entry
     // because the non-ROCr custom code object parsing is called before
     // iterating over the code object symbols using ROCr
     if (KernelInfoTable[gpu].find(kernelName) == KernelInfoTable[gpu].end()) {
       return HSA_STATUS_ERROR;
     }
     // found, so assign and update
     info = KernelInfoTable[gpu][kernelName];