Remove depr hsa_code_object_t - find ISA from ELF

Required to use new hsa_code_object_reader_t for previous usage of hsa_code_object_t. Remove deprecated functions hsa_code_object_deserialize and hsa_executable_load_code_object. Also isCompatible ISA should find the target from ELF header to make decision. Remove code_object from ISA verification.
ROCm · May 9, 2018 · e68f24a · e68f24a
1 parent 2883d82
commit e68f24a
Show file tree

Hide file tree

Showing 3 changed files with 57 additions and 26 deletions.
diff --git a/benchmarks/benchEmptyKernel/hsacodelib.CPP b/benchmarks/benchEmptyKernel/hsacodelib.CPP
@@ -12,7 +12,7 @@
 
 
 
-// Load HSACO 
+// Load HSACO
 Kernel load_hsaco(hc::accelerator_view *av, const char * fileName, const char *kernelName)
 {
     hsa_region_t systemRegion = *(hsa_region_t*)av->get_hsa_am_system_region();
@@ -34,17 +34,17 @@ Kernel load_hsaco(hc::accelerator_view *av, const char * fileName, const char *k
 
         hsa_status_t status;
 
-        hsa_code_object_t code_object = {0};
-        status = hsa_code_object_deserialize(&buffer[0], fsize, NULL, &code_object);
+        hsa_code_object_reader_t co_reader = {};
+        status = hsa_code_object_reader_create_from_memory(&buffer[0], fsize, &co_reader);
         assert(status == HSA_STATUS_SUCCESS);
 
         hsa_executable_t hsaExecutable;
         status = hsa_executable_create_alt(HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT,
                                        NULL, &hsaExecutable);
         assert(status == HSA_STATUS_SUCCESS);
 
-        // Load the code object.
-        status = hsa_executable_load_code_object(hsaExecutable, hsaAgent, code_object, NULL);
+        // Load the agent code object.
+        status = hsa_executable_load_agent_code_object(hsaExecutable, hsaAgent, co_reader, NULL, NULL);
         assert(status == HSA_STATUS_SUCCESS);
 
         // Freeze the executable.
@@ -82,7 +82,7 @@ Kernel load_hsaco(hc::accelerator_view *av, const char * fileName, const char *k
 }
 
 
-// Dispatch a GL packet - 
+// Dispatch a GL packet -
 // Convert to AQL and call dispatch_hsa_kernel
 void dispatch_glp_kernel(const grid_launch_parm *lp, const Kernel &k, void *args, int argSize, bool systemScope)
 {
@@ -106,7 +106,7 @@ void dispatch_glp_kernel(const grid_launch_parm *lp, const Kernel &k, void *args
     aql.header =   (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE);
 
     //aql.header |=   (1 << HSA_PACKET_HEADER_BARRIER);
- 
+
     if (systemScope) {
       aql.header |= (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
                     (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);

diff --git a/lib/hsa/mcwamp_hsa.cpp b/lib/hsa/mcwamp_hsa.cpp
@@ -2666,7 +2666,30 @@ class HSADevice final : public KalmarDevice
         }
     }
 
+    inline
+    std::string get_isa_name_from_triple(std::string triple)
+    {
+        static hsa_isa_t tmp{};
+        static const bool is_old_rocr{
+            hsa_isa_from_name(triple.c_str(), &tmp) != HSA_STATUS_SUCCESS};
+
+        if (is_old_rocr) {
+            auto tmp {triple.substr(triple.rfind('x') + 1)};
+            triple.replace(0, std::string::npos, "AMD:AMDGPU");
+
+            for (auto&& x : tmp) {
+                triple.push_back(':');
+                triple.push_back(x);
+            }
+        }
+
+        return triple;
+    }
+
     bool IsCompatibleKernel(void* size, void* source) override {
+        using namespace ELFIO;
+        using namespace std;
+
         hsa_status_t status;
 
         // Allocate memory for kernel source
@@ -2675,24 +2698,32 @@ class HSADevice final : public KalmarDevice
         memcpy(kernel_source, source, kernel_size);
         kernel_source[kernel_size] = '\0';
 
-        // Deserialize code object.
-        hsa_code_object_t code_object = {0};
-        status = hsa_code_object_deserialize(kernel_source, kernel_size, NULL, &code_object);
-        STATUS_CHECK(status, __LINE__);
-        assert(0 != code_object.handle);
+        // Set up ELF header reader
+        elfio reader;
+        istringstream kern_stream{string{
+            kernel_source,
+            kernel_source + kernel_size}};
+        reader.load(kern_stream);
+
+        // Get ISA from ELF header
+        std::string triple = "amdgcn-amd-amdhsa--gfx";
+        unsigned MACH = reader.get_flags() & 0x000000ff;
+
+        switch(MACH) {
+            case 0x023 : triple.append("701"); break;
+            case 0x02a : triple.append("803"); break;
+            case 0x02c : triple.append("900"); break;
+        }
+
+        const auto isa{get_isa_name_from_triple(std::move(triple))};
 
-        // Get ISA of the code object
-        hsa_isa_t code_object_isa;
-        status = hsa_code_object_get_info(code_object, HSA_CODE_OBJECT_INFO_ISA, &code_object_isa);
+        hsa_isa_t co_isa{};
+        status = hsa_isa_from_name(isa.c_str(), &co_isa);
         STATUS_CHECK(status, __LINE__);
 
         // Check if the code object is compatible with ISA of the agent
         bool isCompatible = false;
-        status = hsa_isa_compatible(code_object_isa, agentISA, &isCompatible);
-        STATUS_CHECK(status, __LINE__);
-
-        // Destroy code object
-        status = hsa_code_object_destroy(code_object);
+        status = hsa_isa_compatible(co_isa, agentISA, &isCompatible);
         STATUS_CHECK(status, __LINE__);
 
         // release allocated memory

diff --git a/tests/Unit/DispatchAql/hsacodelib.CPP b/tests/Unit/DispatchAql/hsacodelib.CPP
@@ -12,7 +12,7 @@
 
 
 
-// Load HSACO 
+// Load HSACO
 Kernel load_hsaco(hc::accelerator_view *av, const char * fileName, const char *kernelName)
 {
     hsa_region_t systemRegion = *(hsa_region_t*)av->get_hsa_am_system_region();
@@ -33,17 +33,17 @@ Kernel load_hsaco(hc::accelerator_view *av, const char * fileName, const char *k
 
         hsa_status_t status;
 
-        hsa_code_object_t code_object = {0};
-        status = hsa_code_object_deserialize(&buffer[0], fsize, NULL, &code_object);
+        hsa_code_object_reader_t co_reader = {};
+        status = hsa_code_object_reader_create_from_memory(&buffer[0], fsize, &co_reader);
         assert(status == HSA_STATUS_SUCCESS);
 
         hsa_executable_t hsaExecutable;
         status = hsa_executable_create_alt(HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT,
                                        NULL, &hsaExecutable);
         assert(status == HSA_STATUS_SUCCESS);
 
-        // Load the code object.
-        status = hsa_executable_load_code_object(hsaExecutable, hsaAgent, code_object, NULL);
+        // Load the agent code object.
+        status = hsa_executable_load_agent_code_object(hsaExecutable, hsaAgent, co_reader, NULL, NULL);
         assert(status == HSA_STATUS_SUCCESS);
 
         // Freeze the executable.
@@ -81,7 +81,7 @@ Kernel load_hsaco(hc::accelerator_view *av, const char * fileName, const char *k
 }
 
 
-// Dispatch a GL packet - 
+// Dispatch a GL packet -
 // Convert to AQL and call dispatch_hsa_kernel
 void dispatch_glp_kernel(const grid_launch_parm *lp, const Kernel &k, void *args, int argSize)
 {