SWDEV-198415 - Implement Target ID Proposal

Changes from Jatin Chaudhary, Tony Tye

Change-Id: I14904920ad2356ae2098cf827b83613089d33962

bin/hipcc

@@ -1,5 +1,23 @@
 #!/usr/bin/perl -w
-
+# Copyright (c) 2015-2020 - present Advanced Micro Devices, Inc. All rights reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+# copies of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in
+# all copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+# THE SOFTWARE.
 
 # Need perl > 5.10 to use logic-defined or
 use 5.006; use v5.10.1;
@@ -16,7 +34,7 @@ use Cwd 'abs_path';
 # options are appropriate for the target compiler.
 
 # Environment variable HIP_PLATFORM is to detect amd/nvidia path:
-# HIP_PLATFORM='nvidia' or HIP_PLATFORM='amd'.
+# HIP_PLATFORM='nvidia' or HIP_PLATFORM='amd'/'hcc'.
 # If HIP_PLATFORM is not set hipcc will attempt auto-detect based on if nvcc is found.
 #
 # Other environment variable controls:
@@ -63,7 +81,12 @@ $HIPCC_COMPILE_FLAGS_APPEND=$ENV{'HIPCC_COMPILE_FLAGS_APPEND'};
 $HIPCC_LINK_FLAGS_APPEND=$ENV{'HIPCC_LINK_FLAGS_APPEND'};
 
 # Known HIP target names.
-@knownTargets = ('gfx701', 'gfx801', 'gfx802', 'gfx803', 'gfx900', 'gfx906', 'gfx908', 'gfx90a', 'gfx1010', 'gfx1011', 'gfx1012', 'gfx1030', 'gfx1031');
+@knownTargets = ('gfx700', 'gfx701', 'gfx702', 'gfx703', 'gfx704', 'gfx705',
+                 'gfx801', 'gfx802', 'gfx803', 'gfx805', 'gfx810',
+                 'gfx900', 'gfx902', 'gfx904', 'gfx906', 'gfx908', 'gfx909',
+                 'gfx1010', 'gfx1011', 'gfx1012', 'gfx1030', 'gfx1031', 'gfx1032');
+# Known Features
+@knownFeatures = ('sramecc-', 'sramecc+', 'xnack-', 'xnack+');
 
 #
 # TODO: Fix rpath LDFLAGS settings
@@ -122,7 +145,7 @@ sub delete_temp_dirs {
 }
 
 #---
-#HIP_PLATFORM controls whether to use amd or nvidia as the platform:
+#HIP_PLATFORM controls whether to use amd/hcc or nvidia as the platform:
 #HIP_COMPILER controls whether to use hcc, clang or nvcc for compilation:
 #HIP_RUNTIME controls whether to use hcc, rocclr, or nvcc as the runtime:
 if ($isWindows) {
@@ -163,7 +186,7 @@ if (defined $HIP_ROCCLR_HOME) {
     }
 }
 
-if (defined $HIP_COMPILER and $HIP_PLATFORM eq "amd" and $HIP_COMPILER eq "clang") {
+if (defined $HIP_COMPILER and ($HIP_PLATFORM eq "amd" or $HIP_PLATFORM eq "hcc") and $HIP_COMPILER eq "clang") {
   if (!defined $HIP_CLANG_PATH) {
     $HIP_CLANG_PATH = "$ROCM_PATH/llvm/bin";
   }
@@ -191,7 +214,7 @@ $setStdLib = 0;  # TODO - set to 0
 
 $default_amdgpu_target = 1;
 
-if ($HIP_PLATFORM eq "amd" and $HIP_COMPILER eq "clang") {
+if (($HIP_PLATFORM eq "amd" or $HIP_PLATFORM eq "hcc") and $HIP_COMPILER eq "clang") {
     $HIPCC="$HIP_CLANG_PATH/clang++";
 
     # If $HIPCC clang++ is not compiled, use clang instead
@@ -249,7 +272,7 @@ if ($HIP_PLATFORM eq "amd" and $HIP_COMPILER eq "clang") {
       $HIPCFLAGS .= " -D__HIP_ROCclr__";
     }
 
-} elsif ($HIP_PLATFORM eq "amd" and $HIP_COMPILER eq "hcc") {
+} elsif (($HIP_PLATFORM eq "amd" or $HIP_PLATFORM eq "hcc") and $HIP_COMPILER eq "hcc") {
     $HIP_INCLUDE_PATH = "$HIP_PATH/include";
     if (! defined $HIP_LIB_PATH) {
         $HIP_LIB_PATH = "$HIP_PATH/lib";
@@ -351,7 +374,7 @@ my $hasC = 0;          # options contain a c-style file
 my $hasCXX = 0;        # options contain a cpp-style file (NVCC must force recognition as GPU file)
 my $hasCU = 0;         # options contain a cu-style file (HCC must force recognition as GPU file)
 my $hasHIP = 0;        # options contain a hip-style file (HIP-Clang must pass offloading options)
-my $needHipHcc = ($HIP_PLATFORM eq 'amd' and $HIP_COMPILER eq 'hcc');      # set if we need to link hip_hcc.o from src tree. (some builds, ie cmake, provide their own)
+my $needHipHcc = (($HIP_PLATFORM eq 'amd' or $HIP_PLATFORM eq "hcc") and $HIP_COMPILER eq 'hcc');      # set if we need to link hip_hcc.o from src tree. (some builds, ie cmake, provide their own)
 my $printHipVersion = 0;    # print HIP version
 my $printCXXFlags = 0;      # print HIPCXXFLAGS
 my $printLDFlags = 0;       # print HIPLDFLAGS
@@ -359,10 +382,7 @@ my $runCmd = 1;
 my $buildDeps = 0;
 my $linkType = 1;
 my $setLinkType = 0;
-my $coFormatv3 = 1;
-if(defined $HIP_COMPILER and $HIP_COMPILER eq "hcc") {
-    $coFormatv3 = 0;
-}
+my $hsacoVersion = 0;
 my $funcSupp = 0;      # enable function support
 my $rdc = 0;           # whether -fgpu-rdc is on
 
@@ -375,7 +395,7 @@ if ($verbose & 0x4) {
 
 # Handle code object generation
 my $ISACMD="";
-if($HIP_PLATFORM eq "amd" and $HIP_COMPILER eq "hcc"){
+if(($HIP_PLATFORM eq "amd" or $HIP_PLATFORM eq "hcc") and $HIP_COMPILER eq "hcc"){
     $ISACMD .= "$HIP_PATH/bin/lpl ";
     if($ARGV[0] eq "--genco"){
         foreach $isaarg (@ARGV[1..$#ARGV]){
@@ -390,7 +410,7 @@ if($HIP_PLATFORM eq "amd" and $HIP_COMPILER eq "hcc"){
     }
 }
 
-if(($HIP_PLATFORM eq "amd" and $HIP_COMPILER eq "hcc")){
+if((($HIP_PLATFORM eq "amd" or $HIP_PLATFORM eq "hcc") and $HIP_COMPILER eq "hcc")){
     $ENV{HCC_EXTRA_LIBRARIES}="\n";
 }
 
@@ -412,7 +432,11 @@ if($HIP_PLATFORM eq "nvidia"){
 # TODO: convert toolArgs to an array rather than a string
 my $toolArgs = "";  # arguments to pass to the hcc or nvcc tool
 my $optArg = ""; # -O args
-my $targetOpt = '--amdgpu-target=';
+
+# TODO: hipcc uses --amdgpu-target for historical reasons. It should be replaced
+# by clang option --offload-arch.
+my @targetOpts = ('--offload-arch=', '--amdgpu-target=');
+
 my $targetsStr = "";
 my $skipOutputFile = 0; # file followed by -o should not contibute in picking compiler flags
 my $prevArg = ""; # previous argument
@@ -451,31 +475,25 @@ foreach $arg (@ARGV)
         $setStdLib = 1;
     }
 
-    # Check target selection option: --amdgpu-target=...
-    if (substr($arg, 0, length($targetOpt)) eq $targetOpt) {
-         # If targets string is not empty, add a comma before adding new target option value.
-         $targetsStr .= ($targetsStr ? ',' : '');
-         $targetsStr .=  substr($arg, length($targetOpt));
-         $default_amdgpu_target = 0;
-         # hip-clang does not accept --amdgpu-target= options.
-         if ($HIP_PLATFORM eq "amd" and $HIP_COMPILER eq "clang") {
-             $swallowArg = 1;
-         }
+    # Check target selection option: --offload-arch= and --amdgpu-target=...
+    foreach my $targetOpt (@targetOpts) {
+        if (substr($arg, 0, length($targetOpt)) eq $targetOpt) {
+             # If targets string is not empty, add a comma before adding new target option value.
+             $targetsStr .= ($targetsStr ? ',' : '');
+             $targetsStr .=  substr($arg, length($targetOpt));
+             $default_amdgpu_target = 0;
+             # Collect the GPU arch options and pass them to clang later.
+             if (($HIP_PLATFORM eq "amd" or $HIP_PLATFORM eq "hcc") and $HIP_COMPILER eq "clang") {
+                 $swallowArg = 1;
+             }
+        }
     }
 
-    # code object format parsing
-    if ($trimarg eq '-mcode-object-v3') {
-        $coFormatv3 = 1;
-    }
-    if ($trimarg eq '-mno-code-object-v3') {
-        $coFormatv3 = 0;
-    }
-
-    if (($arg =~ /--genco/) and $HIP_PLATFORM eq 'amd' and $HIP_COMPILER eq 'clang' ) {
+    if (($arg =~ /--genco/) and ($HIP_PLATFORM eq 'amd' or $HIP_PLATFORM eq "hcc") and $HIP_COMPILER eq 'clang' ) {
         $arg = "--cuda-device-only";
     }
 
-    if(($trimarg eq '-stdlib=libstdc++') and ($setStdLib eq 0) and $HIP_PLATFORM eq 'amd' and $HIP_COMPILER eq 'hcc')
+    if(($trimarg eq '-stdlib=libstdc++') and ($setStdLib eq 0) and ($HIP_PLATFORM eq 'amd' or $HIP_PLATFORM eq "hcc") and $HIP_COMPILER eq 'hcc')
     {
         $HIPCXXFLAGS .= $HCC_WA_FLAGS;
         $setStdLib = 1;
@@ -518,13 +536,19 @@ foreach $arg (@ARGV)
     {
         $optArg = $arg;
     }
+    if($arg =~ '--amdhsa-code-object-version=')
+    {
+        $arg =~ s/--amdhsa-code-object-version=//;
+        $hsacoVersion = $arg;
+        $swallowArg = 1;
+    }
 
     ## process linker response file for hip-clang
     ## extract object files from static library and pass them directly to
     ## hip-clang in command line.
     ## ToDo: Remove this after hip-clang switch to lto and lld is able to
     ## handle clang-offload-bundler bundles.
-    if ($arg =~ m/^-Wl,@/ and $HIP_PLATFORM eq 'amd' and $HIP_COMPILER eq 'clang') {
+    if ($arg =~ m/^-Wl,@/ and ($HIP_PLATFORM eq 'amd' or $HIP_PLATFORM eq "hcc") and $HIP_COMPILER eq 'clang') {
         my $file = substr $arg, 5;
         open my $in, "<:encoding(utf8)", $file or die "$file: $!";
         my $new_arg = "";
@@ -581,7 +605,7 @@ foreach $arg (@ARGV)
         $arg = "$new_arg -Wl,\@$new_file";
         $escapeArg = 0;
     } elsif (($arg =~ m/\.a$/ || $arg =~ m/\.lo$/) &&
-             $HIP_PLATFORM eq 'amd' and $HIP_COMPILER eq 'clang') {
+             ($HIP_PLATFORM eq 'amd' or $HIP_PLATFORM eq "hcc") and $HIP_COMPILER eq 'clang') {
         ## process static library for hip-clang
         ## extract object files from static library and pass them directly to
         ## hip-clang.
@@ -678,13 +702,13 @@ foreach $arg (@ARGV)
                 $needCXXFLAGS = 1;
                 if ($HIP_COMPILE_CXX_AS_HIP eq '0' or $HIP_COMPILER ne "clang") {
                     $hasCXX = 1;
-                } elsif ($HIP_PLATFORM eq "amd" and $HIP_COMPILER eq "clang") {
+                } elsif (($HIP_PLATFORM eq "amd" or $HIP_PLATFORM eq "hcc") and $HIP_COMPILER eq "clang") {
                     $hasHIP = 1;
                     $toolArgs .= " -x hip";
                 }
             } elsif ((($arg =~ /\.cu$/ or $arg =~ /\.cuh$/) and $HIP_COMPILE_CXX_AS_HIP ne '0') or ($arg =~ /\.hip$/)) {
                 $needCXXFLAGS = 1;
-                if ($HIP_PLATFORM eq "amd" and $HIP_COMPILER eq "clang") {
+                if (($HIP_PLATFORM eq "amd" or $HIP_PLATFORM eq "hcc") and $HIP_COMPILER eq "clang") {
                     $hasHIP = 1;
                     $toolArgs .= " -x hip";
                 } else {
@@ -713,7 +737,7 @@ foreach $arg (@ARGV)
     $prevArg = $arg;
 }
 
-if($HIP_PLATFORM eq "amd"){
+if($HIP_PLATFORM eq "amd" or $HIP_PLATFORM eq "hcc"){
     # No AMDGPU target specified at commandline. So look for HCC_AMDGPU_TARGET
     if($default_amdgpu_target eq 1) {
         if (defined $ENV{HCC_AMDGPU_TARGET}) {
@@ -733,14 +757,25 @@ if($HIP_PLATFORM eq "amd"){
     if($HIP_COMPILER eq "hcc") {
         $GPU_ARCH_OPT = " --amdgpu-target=";
     } else {
-        $GPU_ARCH_OPT = " --cuda-gpu-arch=";
+        $GPU_ARCH_OPT = " --offload-arch=";
     }
 
     foreach my $val (@targets) {
         # Ignore 'gfx000' target reported by rocm_agent_enumerator.
         if ($val ne 'gfx000') {
-            # Construct an arch macro to be passed to the compiler.
-            # Add the arch option and macro to the compiler options.
+            my @procAndFeatures = split(':', $val);
+            $len = scalar @procAndFeatures;
+            my $procName;
+            if($len ge 1 and $len le 3) { # proc and features
+                $procName = $procAndFeatures[0];
+                for my $i (1 .. $#procAndFeatures) {
+                    if (grep($procAndFeatures[$i], @knownFeatures) eq 0) {
+                        print "Warning: The Feature: $procAndFeatures[$i] is unknown. Correct compilation is not guaranteed.\n";
+                    }
+                }
+            } else {
+                $procName = $val;
+            }
             $GPU_ARCH_ARG = $GPU_ARCH_OPT . $val;
             $HIPLDARCHFLAGS .= $GPU_ARCH_ARG;
             if ($HIP_COMPILER eq 'clang' and $hasHIP) {
@@ -748,31 +783,43 @@ if($HIP_PLATFORM eq "amd"){
             }
 
             # If the specified target is not in the list of known target names, emit a warning.
-            if (grep(/$val/, @knownTargets) eq 0) {
+            if (grep($procName, @knownTargets) eq 0) {
                 print "Warning: The specified HIP target: $val is unknown. Correct compilation is not guaranteed.\n";
             }
         }
     }
+    if ($hsacoVersion > 0) {
+        if ($HIP_COMPILER eq "hcc") {
+            if ($hsacoVersion == 2) {
+                $HIPLDFLAGS .= " -mno-code-object-v3";
+                $HIPCXXFLAGS .= " -mno-code-object-v3";
+            } elsif ($hsacoVersion == 3) {
+                $HIPLDFLAGS .= " -mcode-object-v3";
+                $HIPCXXFLAGS .= " -mcode-object-v3";
+            } else {
+                print "hcc does not support --amdhsa-code-object-version=$hsacoVersion. Only code object version 2 and 3 is supported.\n" and die();
+            }
+        } else {
+            if ($compileOnly eq 0) {
+                $HIPLDFLAGS .= " -mcode-object-version=$hsacoVersion";
+            } else {
+                $HIPCXXFLAGS .= " -mcode-object-version=$hsacoVersion";
+            }
+        }
+    }
 
     # rocm_agent_enumerator failed! Throw an error and die if linking is required
     if ($default_amdgpu_target eq 1 and $compileOnly eq 0) {
-        print "No valid AMD GPU target was either specified or found. Please specify a valid target using --amdgpu-target=" and die();
+        print "No valid AMD GPU target was either specified or found. Please specify a valid target using --offload-arch=<target>.\n" and die();
     }
 
     $ENV{HCC_EXTRA_LIBRARIES}="\n";
 }
 
-# hcc defaults to v2, so we need to convert to the appropriate flag
-# hip-clang defaults to v3, so we don't need to do anything
-if ($coFormatv3 and $HIP_PLATFORM eq 'amd' and $HIP_COMPILER eq 'hcc') {
-    $HIPLDFLAGS .= " -mcode-object-v3";
-    $HIPCXXFLAGS .= " -mcode-object-v3";
-}
-
 if ($hasCXX and $HIP_PLATFORM eq 'nvidia') {
     $HIPCXXFLAGS .= " -x cu";
 }
-if ($hasCU and $HIP_PLATFORM eq 'amd' and $HIP_COMPILER eq 'hcc') {
+if ($hasCU and ($HIP_PLATFORM eq 'amd' or $HIP_PLATFORM eq "hcc") and $HIP_COMPILER eq 'hcc') {
     $HIPCXXFLAGS .= " -x c++";
 }
 
@@ -781,17 +828,17 @@ if ($buildDeps and $HIP_PLATFORM eq 'nvidia') {
     $HIPCFLAGS .= " -M -D__CUDACC__";
 }
 
-if ($buildDeps and $HIP_PLATFORM eq 'amd' and $HIP_COMPILER eq 'clang') {
+if ($buildDeps and ($HIP_PLATFORM eq 'amd' or $HIP_PLATFORM eq "hcc") and $HIP_COMPILER eq 'clang') {
     $HIPCXXFLAGS .= " --cuda-host-only";
 }
 
 # Add --hip-link only if it is compile only and -fgpu-rdc is on.
-if ($rdc and !$compileOnly and $HIP_PLATFORM eq 'amd' and $HIP_COMPILER eq 'clang') {
+if ($rdc and !$compileOnly and ($HIP_PLATFORM eq 'amd' or $HIP_PLATFORM eq "hcc") and $HIP_COMPILER eq 'clang') {
     $HIPLDFLAGS .= " --hip-link";
     $HIPLDFLAGS .= $HIPLDARCHFLAGS;
 }
 
-if ($setStdLib eq 0 and $HIP_PLATFORM eq 'amd' and $HIP_COMPILER eq 'hcc')
+if ($setStdLib eq 0 and ($HIP_PLATFORM eq 'amd' or $HIP_PLATFORM eq "hcc") and $HIP_COMPILER eq 'hcc')
 {
     $HIPCXXFLAGS .= $HCC_WA_FLAGS;
 }
@@ -810,7 +857,7 @@ if ($needHipHcc) {
 # Reason is that NVCC uses the file extension to determine whether to compile in CUDA mode or
 # pass-through CPP mode.
 
-if ($HIP_PLATFORM eq "amd" and $HIP_COMPILER eq "clang") {
+if (($HIP_PLATFORM eq "amd" or $HIP_PLATFORM eq "hcc") and $HIP_COMPILER eq "clang") {
     # Set default optimization level to -O3 for hip-clang.
     if ($optArg eq "") {
         $HIPCXXFLAGS .= " -O3";
@@ -888,7 +935,7 @@ if ($printLDFlags) {
     print $HIPLDFLAGS;
 }
 if ($runCmd) {
-    if ($HIP_PLATFORM eq "amd" and $HIP_COMPILER eq "hcc" and exists($hipConfig{'HCC_VERSION'}) and $HCC_VERSION ne $hipConfig{'HCC_VERSION'}) {
+    if (($HIP_PLATFORM eq "amd" or $HIP_PLATFORM eq "hcc") and $HIP_COMPILER eq "hcc" and exists($hipConfig{'HCC_VERSION'}) and $HCC_VERSION ne $hipConfig{'HCC_VERSION'}) {
         print ("HIP ($HIP_PATH) was built using hcc $hipConfig{'HCC_VERSION'}, but you are using $HCC_HOME/hcc with version $HCC_VERSION from hipcc. Please rebuild HIP including cmake or update HCC_HOME variable.\n") ;
         die unless $ENV{'HIP_IGNORE_HCC_VERSION'};
     }

docs/markdown/hip_porting_guide.md

@@ -349,6 +349,13 @@ MY_LAUNCH (hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad), true, "firstCal
 hipcc is a portable compiler driver that will call nvcc or HIP-Clang (depending on the target system) and attach all required include and library options. It passes options through to the target compiler. Tools that call hipcc must ensure the compiler options are appropriate for the target compiler.
 The `hipconfig` script may helpful in identifying the target platform, compiler and runtime. It can also help set options appropriately.
 
+### Option for specifying GPU processor
+
+`--offload-arch=X` can be used to specify [target ID](https://clang.llvm.org/docs/ClangOffloadBundlerFileFormat.html#target-id).
+
+For backward compatibility, hipcc also accepts `--amdgpu-target=X` for specifying target ID. However, it will be deprecated
+in future releases.
+
 ## Linking Issues
 
 ### Linking With hipcc

hip-config.cmake.in

@@ -183,10 +183,10 @@ if(HIP_COMPILER STREQUAL "clang")
 
   foreach(GPU_TARGET ${GPU_TARGETS})
       set_property(TARGET hip::device APPEND PROPERTY
-        INTERFACE_COMPILE_OPTIONS "--cuda-gpu-arch=${GPU_TARGET}"
+        INTERFACE_COMPILE_OPTIONS "--offload-arch=${GPU_TARGET}"
       )
       set_property(TARGET hip::device APPEND PROPERTY
-        INTERFACE_LINK_LIBRARIES "--cuda-gpu-arch=${GPU_TARGET}"
+        INTERFACE_LINK_LIBRARIES "--offload-arch=${GPU_TARGET}"
       )
   endforeach()
   #Add support for parallel build and link

rocclr/hip_code_object.cpp

@@ -1,3 +1,25 @@
+/*
+Copyright (c) 2015-2020 - present Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
 #include "hip_code_object.hpp"
 
 #include <cstring>
@@ -8,25 +30,312 @@
 #include "platform/program.hpp"
 #include <elf/elf.hpp>
 
+namespace {
+size_t constexpr strLiteralLength(char const* str) {
+  return *str ? 1 + strLiteralLength(str + 1) : 0;
+}
+constexpr char const* CLANG_OFFLOAD_BUNDLER_MAGIC_STR = "__CLANG_OFFLOAD_BUNDLE__";
+constexpr char const* OFFLOAD_KIND_HIP = "hip";
+constexpr char const* OFFLOAD_KIND_HIPV4 = "hipv4";
+constexpr char const* OFFLOAD_KIND_HCC = "hcc";
+constexpr char const* AMDGCN_TARGET_TRIPLE = "amdgcn-amd-amdhsa-";
+
+// ClangOFFLOADBundle info.
+static constexpr size_t bundle_magic_string_size =
+    strLiteralLength(CLANG_OFFLOAD_BUNDLER_MAGIC_STR);
+
+// Clang Offload bundler description & Header.
+struct __ClangOffloadBundleInfo {
+  uint64_t offset;
+  uint64_t size;
+  uint64_t bundleEntryIdSize;
+  const char bundleEntryId[1];
+};
+
+struct __ClangOffloadBundleHeader {
+  const char magic[bundle_magic_string_size - 1];
+  uint64_t numOfCodeObjects;
+  __ClangOffloadBundleInfo desc[1];
+};
+}  // namespace
+
 namespace hip {
 
 uint64_t CodeObject::ElfSize(const void *emi) {
   return amd::Elf::getElfSize(emi);
 }
 
-bool CodeObject::isCompatibleCodeObject(const std::string& codeobj_target_id,
-                                        const char* device_name) {
-  // Workaround for device name mismatch.
-  // Device name may contain feature strings delimited by '+', e.g.
-  // gfx900+xnack. Currently HIP-Clang does not include feature strings
-  // in code object target id in fat binary. Therefore drop the feature
-  // strings from device name before comparing it with code object target id.
-  std::string short_name(device_name);
-  auto feature_loc = short_name.find('+');
-  if (feature_loc != std::string::npos) {
-    short_name.erase(feature_loc);
+static bool getProcName(uint32_t EFlags, std::string& proc_name, bool& xnackSupported,
+                        bool& sramEccSupported) {
+  switch (EFlags & EF_AMDGPU_MACH) {
+    case EF_AMDGPU_MACH_AMDGCN_GFX700:
+      xnackSupported = false;
+      sramEccSupported = false;
+      proc_name = "gfx700";
+      break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX701:
+      xnackSupported = false;
+      sramEccSupported = false;
+      proc_name = "gfx701";
+      break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX702:
+      xnackSupported = false;
+      sramEccSupported = false;
+      proc_name = "gfx702";
+      break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX703:
+      xnackSupported = false;
+      sramEccSupported = false;
+      proc_name = "gfx703";
+      break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX704:
+      xnackSupported = false;
+      sramEccSupported = false;
+      proc_name = "gfx704";
+      break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX801:
+      xnackSupported = true;
+      sramEccSupported = false;
+      proc_name = "gfx801";
+      break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX802:
+      xnackSupported = false;
+      sramEccSupported = false;
+      proc_name = "gfx802";
+      break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX803:
+      xnackSupported = false;
+      sramEccSupported = false;
+      proc_name = "gfx803";
+      break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX810:
+      xnackSupported = true;
+      sramEccSupported = false;
+      proc_name = "gfx810";
+      break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX900:
+      xnackSupported = true;
+      sramEccSupported = false;
+      proc_name = "gfx900";
+      break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX902:
+      xnackSupported = true;
+      sramEccSupported = false;
+      proc_name = "gfx902";
+      break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX904:
+      xnackSupported = true;
+      sramEccSupported = false;
+      proc_name = "gfx904";
+      break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX906:
+      xnackSupported = true;
+      sramEccSupported = true;
+      proc_name = "gfx906";
+      break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX908:
+      xnackSupported = true;
+      sramEccSupported = true;
+      proc_name = "gfx908";
+      break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX909:
+      xnackSupported = true;
+      sramEccSupported = false;
+      proc_name = "gfx909";
+      break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX1010:
+      xnackSupported = true;
+      sramEccSupported = false;
+      proc_name = "gfx1010";
+      break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX1011:
+      xnackSupported = true;
+      sramEccSupported = false;
+      proc_name = "gfx1011";
+      break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX1012:
+      xnackSupported = true;
+      sramEccSupported = false;
+      proc_name = "gfx1012";
+      break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX1030:
+      xnackSupported = false;
+      sramEccSupported = false;
+      proc_name = "gfx1030";
+      break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX1031:
+      xnackSupported = false;
+      sramEccSupported = false;
+      proc_name = "gfx1031";
+      break;
+    default:
+      return false;
   }
-  return codeobj_target_id == short_name;
+  return true;
+}
+
+static bool getTripleTargetIDFromCodeObject(const void* code_object, std::string& target_id,
+                                            unsigned& co_version) {
+  if (!code_object) return false;
+  const Elf64_Ehdr* ehdr = reinterpret_cast<const Elf64_Ehdr*>(code_object);
+  if (ehdr->e_machine != EM_AMDGPU) return false;
+  if (ehdr->e_ident[EI_OSABI] != ELFOSABI_AMDGPU_HSA) return false;
+
+  bool isXnackSupported{false}, isSramEccSupported{false};
+
+  std::string proc_name;
+  if (!getProcName(ehdr->e_flags, proc_name, isXnackSupported, isSramEccSupported)) return false;
+  target_id = std::string(AMDGCN_TARGET_TRIPLE) + '-' + proc_name;
+
+  switch (ehdr->e_ident[EI_ABIVERSION]) {
+    case ELFABIVERSION_AMDGPU_HSA_V2: {
+      co_version = 2;
+      return false;
+    }
+
+    case ELFABIVERSION_AMDGPU_HSA_V3: {
+      co_version = 3;
+      if (isSramEccSupported) {
+        if (ehdr->e_flags & EF_AMDGPU_FEATURE_SRAMECC_V3)
+          target_id += ":sramecc+";
+        else
+          target_id += ":sramecc-";
+      }
+      if (isXnackSupported) {
+        if (ehdr->e_flags & EF_AMDGPU_FEATURE_XNACK_V3)
+          target_id += ":xnack+";
+        else
+          target_id += ":xnack-";
+      }
+      break;
+    }
+
+    case ELFABIVERSION_AMDGPU_HSA_V4: {
+      co_version = 4;
+      unsigned co_sram_value = (ehdr->e_flags) & EF_AMDGPU_FEATURE_SRAMECC_V4;
+      if (co_sram_value == EF_AMDGPU_FEATURE_SRAMECC_OFF_V4)
+        target_id += ":sramecc-";
+      else if (co_sram_value == EF_AMDGPU_FEATURE_SRAMECC_ON_V4)
+        target_id += ":sramecc+";
+
+      unsigned co_xnack_value = (ehdr->e_flags) & EF_AMDGPU_FEATURE_XNACK_V4;
+      if (co_xnack_value == EF_AMDGPU_FEATURE_XNACK_OFF_V4)
+        target_id += ":xnack-";
+      else if (co_xnack_value == EF_AMDGPU_FEATURE_XNACK_ON_V4)
+        target_id += ":xnack+";
+      break;
+    }
+
+    default: {
+      return false;
+    }
+  }
+  return true;
+}
+
+// Consumes the string 'consume_' from the starting of the given input
+// eg: input = amdgcn-amd-amdhsa--gfx908 and consume_ is amdgcn-amd-amdhsa--
+// input will become gfx908.
+static bool consume(std::string& input, std::string consume_) {
+  if (input.substr(0, consume_.size()) != consume_) {
+    return false;
+  }
+  input = input.substr(consume_.size());
+  return true;
+}
+
+// Trim String till character, will be used to get gpuname
+// example: input is gfx908:sram-ecc+ and trim char is :
+// input will become sram-ecc+.
+static std::string trimName(std::string& input, char trim) {
+  auto pos_ = input.find(trim);
+  auto res = input;
+  if (pos_ == std::string::npos) {
+    input = "";
+  } else {
+    res = input.substr(0, pos_);
+    input = input.substr(pos_);
+  }
+  return res;
+}
+
+static char getFeatureValue(std::string& input, std::string feature) {
+  char res = ' ';
+  if (consume(input, std::move(feature))) {
+    res = input[0];
+    input = input.substr(1);
+  }
+  return res;
+}
+
+static bool getTargetIDValue(std::string& input, std::string& processor, char& sramecc_value,
+                             char& xnack_value) {
+  processor = trimName(input, ':');
+  sramecc_value = getFeatureValue(input, std::string(":sramecc"));
+  if (sramecc_value != ' ' && sramecc_value != '+' && sramecc_value != '-') return false;
+  xnack_value = getFeatureValue(input, std::string(":xnack"));
+  if (xnack_value != ' ' && xnack_value != '+' && xnack_value != '-') return false;
+  return true;
+}
+
+static bool getTripleTargetID(std::string bundled_co_entry_id, const void* code_object,
+                              std::string& co_triple_target_id, unsigned& co_version) {
+  std::string offload_kind = trimName(bundled_co_entry_id, '-');
+  if (offload_kind != OFFLOAD_KIND_HIPV4 && offload_kind != OFFLOAD_KIND_HIP &&
+      offload_kind != OFFLOAD_KIND_HCC)
+    return false;
+
+  if (offload_kind != OFFLOAD_KIND_HIPV4)
+    return getTripleTargetIDFromCodeObject(code_object, co_triple_target_id, co_version);
+
+  // For code object V4 onwards the bundled code object entry ID correctly
+  // specifies the target tripple.
+  co_version = 4;
+  co_triple_target_id = bundled_co_entry_id.substr(1);
+  return true;
+}
+
+static bool isCodeObjectCompatibleWithDevice(std::string co_triple_target_id,
+                                             std::string agent_triple_target_id) {
+  // Primitive Check
+  if (co_triple_target_id == agent_triple_target_id) return true;
+
+  // Parse code object triple target id
+  if (!consume(co_triple_target_id, std::string(AMDGCN_TARGET_TRIPLE) + '-')) {
+    return false;
+  }
+
+  std::string co_processor;
+  char co_sram_ecc, co_xnack;
+  if (!getTargetIDValue(co_triple_target_id, co_processor, co_sram_ecc, co_xnack)) {
+    return false;
+  }
+
+  if (!co_triple_target_id.empty()) return false;
+
+  // Parse agent isa triple target id
+  if (!consume(agent_triple_target_id, std::string(AMDGCN_TARGET_TRIPLE) + '-')) {
+    return false;
+  }
+
+  std::string agent_isa_processor;
+  char isa_sram_ecc, isa_xnack;
+  if (!getTargetIDValue(agent_triple_target_id, agent_isa_processor, isa_sram_ecc, isa_xnack)) {
+    return false;
+  }
+
+  if (!agent_triple_target_id.empty()) return false;
+
+  // Check for compatibility
+  if (agent_isa_processor != co_processor) return false;
+  if (co_sram_ecc != ' ') {
+    if (co_sram_ecc != isa_sram_ecc) return false;
+  }
+  if (co_xnack != ' ') {
+    if (co_xnack != isa_xnack) return false;
+  }
+
+  return true;
 }
 
 // This will be moved to COMGR eventually
@@ -72,52 +381,80 @@ hipError_t CodeObject::ExtractCodeObjectFromMemory(const void* data,
 }
 
 // This will be moved to COMGR eventually
-hipError_t CodeObject::extractCodeObjectFromFatBinary(const void* data,
-                       const std::vector<const char*>& device_names,
+hipError_t CodeObject::extractCodeObjectFromFatBinary(const void* data, 
+                       const std::vector<const char*>& agent_triple_target_ids,
                        std::vector<std::pair<const void*, size_t>>& code_objs) {
-  std::string magic((const char*)data, sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC_STR) - 1);
+  std::string magic((const char*)data, bundle_magic_string_size);
   if (magic.compare(CLANG_OFFLOAD_BUNDLER_MAGIC_STR)) {
     return hipErrorInvalidKernelFile;
   }
 
-  code_objs.resize(device_names.size());
+  // Initialize Code objects
+  code_objs.reserve(agent_triple_target_ids.size());
+  for (size_t i = 0; i < agent_triple_target_ids.size(); i++) {
+    code_objs.push_back(std::make_pair(nullptr, 0));
+  }
+
   const auto obheader = reinterpret_cast<const __ClangOffloadBundleHeader*>(data);
   const auto* desc = &obheader->desc[0];
-  unsigned num_code_objs = 0;
-  for (uint64_t i = 0; i < obheader->numBundles; ++i,
-       desc = reinterpret_cast<const __ClangOffloadBundleDesc*>(
-           reinterpret_cast<uintptr_t>(&desc->triple[0]) + desc->tripleSize)) {
+  size_t num_code_objs = code_objs.size();
+  for (uint64_t i = 0; i < obheader->numOfCodeObjects; ++i,
+                desc = reinterpret_cast<const __ClangOffloadBundleInfo*>(
+                    reinterpret_cast<uintptr_t>(&desc->bundleEntryId[0]) +
+                    desc->bundleEntryIdSize)) {
+    const void* image =
+        reinterpret_cast<const void*>(reinterpret_cast<uintptr_t>(obheader) + desc->offset);
+    const size_t image_size = desc->size;
 
-    std::size_t offset = 0;
-    if (!std::strncmp(desc->triple, HIP_AMDGCN_AMDHSA_TRIPLE,
-        sizeof(HIP_AMDGCN_AMDHSA_TRIPLE) - 1)) {
-      offset = sizeof(HIP_AMDGCN_AMDHSA_TRIPLE); //For code objects created by CLang
-    } else if (!std::strncmp(desc->triple, HCC_AMDGCN_AMDHSA_TRIPLE,
-               sizeof(HCC_AMDGCN_AMDHSA_TRIPLE) - 1)) {
-      offset = sizeof(HCC_AMDGCN_AMDHSA_TRIPLE); //For code objects created by Hcc
-    } else {
-      continue;
-    }
-    std::string target(desc->triple + offset, desc->tripleSize - offset);
+    if (num_code_objs == 0) break;
+    std::string bundleEntryId{desc->bundleEntryId, desc->bundleEntryIdSize};
 
-    const void *image = reinterpret_cast<const void*>(
-        reinterpret_cast<uintptr_t>(obheader) + desc->offset);
-    size_t size = desc->size;
+    unsigned co_version = 0;
+    std::string co_triple_target_id;
+    if (!getTripleTargetID(bundleEntryId, image, co_triple_target_id, co_version)) continue;
 
-    for (size_t dev = 0; dev < device_names.size(); ++dev) {
-      const char* name = device_names[dev];
-
-      if (!isCompatibleCodeObject(target, name)) {
-          continue;
+    for (size_t dev = 0; dev < agent_triple_target_ids.size(); ++dev) {
+      if (code_objs[dev].first) continue;
+      if (isCodeObjectCompatibleWithDevice(co_triple_target_id, agent_triple_target_ids[dev])) {
+        code_objs[dev] = std::make_pair(image, image_size);
+        --num_code_objs;
       }
-      code_objs[dev] = std::make_pair(image, size);
-      num_code_objs++;
     }
   }
-  if (num_code_objs == device_names.size()) {
+  if (num_code_objs == 0) {
     return hipSuccess;
   } else {
-    guarantee(false && "hipErrorNoBinaryForGpu: Coudn't find binary for current devices!");
+    LogPrintfError("%s",
+                   "hipErrorNoBinaryForGpu: Unable to find code object for all current devices!");
+    LogPrintfError("%s", "  Devices:");
+    for (size_t i = 0; i < agent_triple_target_ids.size(); i++) {
+      LogPrintfError("    %s - [%s]", agent_triple_target_ids[i],
+                     ((code_objs[i].first) ? "Found" : "Not Found"));
+    }
+    const auto obheader = reinterpret_cast<const __ClangOffloadBundleHeader*>(data);
+    const auto* desc = &obheader->desc[0];
+    LogPrintfError("%s", "  Bundled Code Objects:");
+    for (uint64_t i = 0; i < obheader->numOfCodeObjects; ++i,
+                  desc = reinterpret_cast<const __ClangOffloadBundleInfo*>(
+                      reinterpret_cast<uintptr_t>(&desc->bundleEntryId[0]) +
+                      desc->bundleEntryIdSize)) {
+      std::string bundleEntryId{desc->bundleEntryId, desc->bundleEntryIdSize};
+      const void* image =
+          reinterpret_cast<const void*>(reinterpret_cast<uintptr_t>(obheader) + desc->offset);
+
+      unsigned co_version = 0;
+      std::string co_triple_target_id;
+      bool valid_co = getTripleTargetID(bundleEntryId, image, co_triple_target_id, co_version);
+
+      if (valid_co) {
+        LogPrintfError("    %s - [code object v%u is %s]", bundleEntryId.c_str(), co_version,
+                       co_triple_target_id.c_str());
+      } else {
+        LogPrintfError("    %s - [Unsupported]", bundleEntryId.c_str());
+      }
+    }
+
+    guarantee(false && "hipErrorNoBinaryForGpu: Unable to find code object for all current devices!");
     return hipErrorNoBinaryForGpu;
   }
 }

rocclr/hip_code_object.hpp

@@ -1,8 +1,31 @@
+/*
+Copyright (c) 2015-2020 - present Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
 #ifndef HIP_CODE_OBJECT_HPP
 #define HIP_CODE_OBJECT_HPP
 
 #include "hip_global.hpp"
 
+#include <cstring>
 #include <unordered_map>
 
 #include "hip/hip_runtime.h"
@@ -18,28 +41,13 @@ namespace hip {
 
 //Code Object base class
 class CodeObject {
-public:
+ public:
   virtual ~CodeObject() {}
 
-  //ClangOFFLOADBundle info
-  #define CLANG_OFFLOAD_BUNDLER_MAGIC_STR "__CLANG_OFFLOAD_BUNDLE__"
-  #define HIP_AMDGCN_AMDHSA_TRIPLE "hip-amdgcn-amd-amdhsa"
-  #define HCC_AMDGCN_AMDHSA_TRIPLE "hcc-amdgcn-amd-amdhsa-"
-
-  //Clang Offload bundler description & Header
-  struct __ClangOffloadBundleDesc {
-    uint64_t offset;
-    uint64_t size;
-    uint64_t tripleSize;
-    const char triple[1];
-  };
-
-  struct __ClangOffloadBundleHeader {
-    const char magic[sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC_STR) - 1];
-    uint64_t numBundles;
-    __ClangOffloadBundleDesc desc[1];
-  };
-
+  // Functions to add_dev_prog and build
+  static hipError_t add_program(int deviceId, hipModule_t hmod, const void* binary_ptr,
+                                size_t binary_size);
+  static hipError_t build_module(hipModule_t hmod, const std::vector<amd::Device*>& devices);
 
   // Given an file desc and file size, extracts to code object for corresponding devices,
   // return code_objs{binary_ptr, binary_size}, which could be used to determine foffset
@@ -57,15 +65,14 @@ public:
   static uint64_t ElfSize(const void* emi);
 
 protected:
+  //Given an ptr to image or file, extracts to code object
+  //for corresponding devices
   static hipError_t extractCodeObjectFromFatBinary(const void*,
                     const std::vector<const char*>&,
                     std::vector<std::pair<const void*, size_t>>&);
 
   CodeObject() {}
 private:
-  static bool isCompatibleCodeObject(const std::string& codeobj_target_id,
-                                     const char* device_name);
-
   friend const std::vector<hipModule_t>& modules();
 };
 
@@ -137,6 +144,6 @@ private:
   std::unordered_map<const void*, Var*> vars_;
 };
 
-}; //namespace: hip
+}; // namespace hip
 
 #endif /* HIP_CODE_OBJECT_HPP */

rocclr/hip_fatbin.cpp

@@ -49,12 +49,12 @@ FatBinaryInfo::~FatBinaryInfo() {
 hipError_t FatBinaryInfo::ExtractFatBinary(const std::vector<hip::Device*>& devices) {
   hipError_t hip_error = hipSuccess;
   std::vector<std::pair<const void*, size_t>> code_objs;
-  code_objs.resize(devices.size());
 
   // Copy device names for Extract Code object File
   std::vector<const char*> device_names;
+  device_names.reserve(devices.size());
   for (size_t dev_idx = 0; dev_idx < devices.size(); ++dev_idx) {
-    device_names.push_back(devices[dev_idx]->devices()[0]->info().name_);
+    device_names.push_back(devices[dev_idx]->devices()[0]->info().targetId_);
   }
 
   // We are given file name, get the file desc and file size
@@ -85,7 +85,7 @@ hipError_t FatBinaryInfo::ExtractFatBinary(const std::vector<hip::Device*>& devi
   }
 
   if (hip_error == hipErrorNoBinaryForGpu) {
-    guarantee(false && "hipErrorNoBinaryForGpu: Coudn't find binary for current devices!");
+    guarantee(false && "hipErrorNoBinaryForGpu: Couldn't find binary for current devices!");
     return hip_error;
   }