Merge branch 'amd-master-next' into amd-npi-next

Change-Id: I4181bbab4fc0a5ec857fe39d6148ee8b4975836d


[ROCm/hip commit: 573b2dced3]

projects/hip/docs/doxygen-input/doxy.cfg

@@ -1,4 +1,4 @@
-# Doxyfile 1.8.9.1
+# Doxyfile 1.8.17
 
 # This file describes the settings to be used by the documentation system
 # doxygen (www.doxygen.org) for a project.
@@ -17,11 +17,11 @@
 # Project related configuration options
 #---------------------------------------------------------------------------
 
-# This tag specifies the encoding used for all characters in the config file
-# that follow. The default is UTF-8 which is also the encoding used for all text
-# before the first occurrence of this tag. Doxygen uses libiconv (or the iconv
-# built into libc) for the transcoding. See http://www.gnu.org/software/libiconv
-# for the list of possible encodings.
+# This tag specifies the encoding used for all characters in the configuration
+# file that follow. The default is UTF-8 which is also the encoding used for all
+# text before the first occurrence of this tag. Doxygen uses libiconv (or the
+# iconv built into libc) for the transcoding. See
+# https://www.gnu.org/software/libiconv/ for the list of possible encodings.
 # The default value is: UTF-8.
 
 DOXYFILE_ENCODING      = UTF-8
@@ -93,6 +93,14 @@ ALLOW_UNICODE_NAMES    = NO
 
 OUTPUT_LANGUAGE        = English
 
+# The OUTPUT_TEXT_DIRECTION tag is used to specify the direction in which all
+# documentation generated by doxygen is written. Doxygen will use this
+# information to generate all generated output in the proper direction.
+# Possible values are: None, LTR, RTL and Context.
+# The default value is: None.
+
+OUTPUT_TEXT_DIRECTION  = None
+
 # If the BRIEF_MEMBER_DESC tag is set to YES, doxygen will include brief member
 # descriptions after the members that are listed in the file and class
 # documentation (similar to Javadoc). Set to NO to disable this.
@@ -118,7 +126,17 @@ REPEAT_BRIEF           = YES
 # the entity):The $name class, The $name widget, The $name file, is, provides,
 # specifies, contains, represents, a, an and the.
 
-ABBREVIATE_BRIEF       =
+ABBREVIATE_BRIEF       = "The $name class" \
+                         "The $name widget" \
+                         "The $name file" \
+                         is \
+                         provides \
+                         specifies \
+                         contains \
+                         represents \
+                         a \
+                         an \
+                         the
 
 # If the ALWAYS_DETAILED_SEC and REPEAT_BRIEF tags are both set to YES then
 # doxygen will generate a detailed section even if there is only a brief
@@ -179,6 +197,16 @@ SHORT_NAMES            = NO
 
 JAVADOC_AUTOBRIEF      = NO
 
+# If the JAVADOC_BANNER tag is set to YES then doxygen will interpret a line
+# such as
+# /***************
+# as being the beginning of a Javadoc-style comment "banner". If set to NO, the
+# Javadoc-style will behave just like regular comments and it will not be
+# interpreted by doxygen.
+# The default value is: NO.
+
+JAVADOC_BANNER         = NO
+
 # If the QT_AUTOBRIEF tag is set to YES then doxygen will interpret the first
 # line (until the first dot) of a Qt-style comment as the brief description. If
 # set to NO, the Qt-style will behave just like regular Qt-style comments (thus
@@ -226,7 +254,12 @@ TAB_SIZE               = 4
 # will allow you to put the command \sideeffect (or @sideeffect) in the
 # documentation, which will result in a user-defined paragraph with heading
 # "Side Effects:". You can put \n's in the value part of an alias to insert
-# newlines.
+# newlines (in the resulting output). You can put ^^ in the value part of an
+# alias to insert a newline as if a physical newline was in the original file.
+# When you need a literal { or } or , in the value part of an alias you have to
+# escape them by means of a backslash (\), this can lead to conflicts with the
+# commands \{ and \} for these it is advised to use the version @{ and @} or use
+# a double escape (\\{ and \\})
 
 ALIASES                =
 
@@ -264,28 +297,37 @@ OPTIMIZE_FOR_FORTRAN   = NO
 
 OPTIMIZE_OUTPUT_VHDL   = NO
 
+# Set the OPTIMIZE_OUTPUT_SLICE tag to YES if your project consists of Slice
+# sources only. Doxygen will then generate output that is more tailored for that
+# language. For instance, namespaces will be presented as modules, types will be
+# separated into more groups, etc.
+# The default value is: NO.
+
+OPTIMIZE_OUTPUT_SLICE  = NO
+
 # Doxygen selects the parser to use depending on the extension of the files it
 # parses. With this tag you can assign which parser to use for a given
 # extension. Doxygen has a built-in mapping, but you can override or extend it
 # using this tag. The format is ext=language, where ext is a file extension, and
-# language is one of the parsers supported by doxygen: IDL, Java, Javascript,
-# C#, C, C++, D, PHP, Objective-C, Python, Fortran (fixed format Fortran:
-# FortranFixed, free formatted Fortran: FortranFree, unknown formatted Fortran:
-# Fortran. In the later case the parser tries to guess whether the code is fixed
-# or free formatted code, this is the default for Fortran type files), VHDL. For
-# instance to make doxygen treat .inc files as Fortran files (default is PHP),
-# and .f files as C (default is Fortran), use: inc=Fortran f=C.
+# language is one of the parsers supported by doxygen: IDL, Java, JavaScript,
+# Csharp (C#), C, C++, D, PHP, md (Markdown), Objective-C, Python, Slice,
+# Fortran (fixed format Fortran: FortranFixed, free formatted Fortran:
+# FortranFree, unknown formatted Fortran: Fortran. In the later case the parser
+# tries to guess whether the code is fixed or free formatted code, this is the
+# default for Fortran type files), VHDL, tcl. For instance to make doxygen treat
+# .inc files as Fortran files (default is PHP), and .f files as C (default is
+# Fortran), use: inc=Fortran f=C.
 #
 # Note: For files without extension you can use no_extension as a placeholder.
 #
 # Note that for custom extensions you also need to set FILE_PATTERNS otherwise
 # the files are not read by doxygen.
 
-EXTENSION_MAPPING += dox=md
+EXTENSION_MAPPING      =
 
 # If the MARKDOWN_SUPPORT tag is enabled then doxygen pre-processes all comments
 # according to the Markdown format, which allows for more readable
-# documentation. See http://daringfireball.net/projects/markdown/ for details.
+# documentation. See https://daringfireball.net/projects/markdown/ for details.
 # The output of markdown processing is further processed by doxygen, so you can
 # mix doxygen, HTML, and XML commands with Markdown formatting. Disable only in
 # case of backward compatibilities issues.
@@ -293,6 +335,15 @@ EXTENSION_MAPPING += dox=md
 
 MARKDOWN_SUPPORT       = YES
 
+# When the TOC_INCLUDE_HEADINGS tag is set to a non-zero value, all headings up
+# to that level are automatically included in the table of contents, even if
+# they do not have an id attribute.
+# Note: This feature currently applies only to Markdown headings.
+# Minimum value: 0, maximum value: 99, default value: 5.
+# This tag requires that the tag MARKDOWN_SUPPORT is set to YES.
+
+TOC_INCLUDE_HEADINGS   = 5
+
 # When enabled doxygen tries to link words that correspond to documented
 # classes, or namespaces to their corresponding documentation. Such a link can
 # be prevented in individual cases by putting a % sign in front of the word or
@@ -318,7 +369,7 @@ BUILTIN_STL_SUPPORT    = NO
 CPP_CLI_SUPPORT        = NO
 
 # Set the SIP_SUPPORT tag to YES if your project consists of sip (see:
-# http://www.riverbankcomputing.co.uk/software/sip/intro) sources only. Doxygen
+# https://www.riverbankcomputing.com/software/sip/intro) sources only. Doxygen
 # will parse them like normal C++ but will assume all classes use public instead
 # of private inheritance when no explicit protection keyword is present.
 # The default value is: NO.
@@ -343,6 +394,13 @@ IDL_PROPERTY_SUPPORT   = YES
 
 DISTRIBUTE_GROUP_DOC   = NO
 
+# If one adds a struct or class to a group and this option is enabled, then also
+# any nested class or struct is added to the same group. By default this option
+# is disabled and one has to add nested compounds explicitly via \ingroup.
+# The default value is: NO.
+
+GROUP_NESTED_COMPOUNDS = NO
+
 # Set the SUBGROUPING tag to YES to allow class member groups of the same type
 # (for instance a group of public functions) to be put as a subgroup of that
 # type (e.g. under the Public Functions section). Set it to NO to prevent
@@ -417,6 +475,12 @@ EXTRACT_ALL            = NO
 
 EXTRACT_PRIVATE        = NO
 
+# If the EXTRACT_PRIV_VIRTUAL tag is set to YES, documented private virtual
+# methods of a class will be included in the documentation.
+# The default value is: NO.
+
+EXTRACT_PRIV_VIRTUAL   = NO
+
 # If the EXTRACT_PACKAGE tag is set to YES, all members with package or internal
 # scope will be included in the documentation.
 # The default value is: NO.
@@ -471,8 +535,8 @@ HIDE_UNDOC_MEMBERS     = NO
 HIDE_UNDOC_CLASSES     = NO
 
 # If the HIDE_FRIEND_COMPOUNDS tag is set to YES, doxygen will hide all friend
-# (class|struct|union) declarations. If set to NO, these declarations will be
-# included in the documentation.
+# declarations. If set to NO, these declarations will be included in the
+# documentation.
 # The default value is: NO.
 
 HIDE_FRIEND_COMPOUNDS  = NO
@@ -495,7 +559,7 @@ INTERNAL_DOCS          = YES
 # names in lower-case letters. If set to YES, upper-case letters are also
 # allowed. This is useful if you have classes or files whose names only differ
 # in case and if your file system supports case sensitive file names. Windows
-# and Mac users are advised to set this option to NO.
+# (including Cygwin) ands Mac users are advised to set this option to NO.
 # The default value is: system dependent.
 
 CASE_SENSE_NAMES       = YES
@@ -609,14 +673,14 @@ GENERATE_TESTLIST      = YES
 # list. This list is created by putting \bug commands in the documentation.
 # The default value is: YES.
 
-GENERATE_BUGLIST       = YES
+GENERATE_BUGLIST       = NO
 
 # The GENERATE_DEPRECATEDLIST tag can be used to enable (YES) or disable (NO)
 # the deprecated list. This list is created by putting \deprecated commands in
 # the documentation.
 # The default value is: YES.
 
-GENERATE_DEPRECATEDLIST= YES
+GENERATE_DEPRECATEDLIST= NO
 
 # The ENABLED_SECTIONS tag can be used to enable conditional documentation
 # sections, marked by \if <section_label> ... \endif and \cond <section_label>
@@ -682,7 +746,7 @@ LAYOUT_FILE            =
 # The CITE_BIB_FILES tag can be used to specify one or more bib files containing
 # the reference definitions. This must be a list of .bib files. The .bib
 # extension is automatically appended if omitted. This requires the bibtex tool
-# to be installed. See also http://en.wikipedia.org/wiki/BibTeX for more info.
+# to be installed. See also https://en.wikipedia.org/wiki/BibTeX for more info.
 # For LaTeX the style of the bibliography can be controlled using
 # LATEX_BIB_STYLE. To use this feature you need bibtex and perl available in the
 # search path. See also \cite for info how to create references.
@@ -727,11 +791,18 @@ WARN_IF_DOC_ERROR      = YES
 # This WARN_NO_PARAMDOC option can be enabled to get warnings for functions that
 # are documented, but have no documentation for their parameters or return
 # value. If set to NO, doxygen will only warn about wrong or incomplete
-# parameter documentation, but not about the absence of documentation.
+# parameter documentation, but not about the absence of documentation. If
+# EXTRACT_ALL is set to YES then this flag will automatically be disabled.
 # The default value is: NO.
 
 WARN_NO_PARAMDOC       = NO
 
+# If the WARN_AS_ERROR tag is set to YES then doxygen will immediately stop when
+# a warning is encountered.
+# The default value is: NO.
+
+WARN_AS_ERROR          = NO
+
 # The WARN_FORMAT tag determines the format of the warning messages that doxygen
 # can produce. The string should contain the $file, $line, and $text tags, which
 # will be replaced by the file and line number from which the warning originated
@@ -755,11 +826,15 @@ WARN_LOGFILE           =
 # The INPUT tag is used to specify the files and/or directories that contain
 # documented source files. You may enter file names like myfile.cpp or
 # directories like /usr/src/myproject. Separate the files or directories with
-# spaces.
+# spaces. See also FILE_PATTERNS and EXTENSION_MAPPING
 # Note: If this tag is empty the current directory is searched.
 
 INPUT                  = $(HIP_PATH)/docs/doxygen-input/mainpage.txt \
+                         $(HIP_PATH)/README.md \
+                         $(HIP_PATH)/CONTRIBUTING.md \
                          $(HIP_PATH)/docs/doxygen-input/sync.txt \
+                         $(HIP_PATH)/INSTALL.md \
+                         $(HIP_PATH)/docs/markdown \
                          $(HIP_PATH)/include/hip \
                          $(HIP_PATH)/include/hip/hcc_detail/ \
                          $(HIP_PATH)/src/
@@ -767,7 +842,7 @@ INPUT                  = $(HIP_PATH)/docs/doxygen-input/mainpage.txt \
 # This tag can be used to specify the character encoding of the source files
 # that doxygen parses. Internally doxygen uses the UTF-8 encoding. Doxygen uses
 # libiconv (or the iconv built into libc) for the transcoding. See the libiconv
-# documentation (see: http://www.gnu.org/software/libiconv) for the list of
+# documentation (see: https://www.gnu.org/software/libiconv/) for the list of
 # possible encodings.
 # The default value is: UTF-8.
 
@@ -775,20 +850,36 @@ INPUT_ENCODING         = UTF-8
 
 # If the value of the INPUT tag contains directories, you can use the
 # FILE_PATTERNS tag to specify one or more wildcard patterns (like *.cpp and
-# *.h) to filter out the source-files in the directories. If left blank the
-# following patterns are tested:*.c, *.cc, *.cxx, *.cpp, *.c++, *.java, *.ii,
-# *.ixx, *.ipp, *.i++, *.inl, *.idl, *.ddl, *.odl, *.h, *.hh, *.hxx, *.hpp,
-# *.h++, *.cs, *.d, *.php, *.php4, *.php5, *.phtml, *.inc, *.m, *.markdown,
-# *.md, *.mm, *.dox, *.py, *.f90, *.f, *.for, *.tcl, *.vhd, *.vhdl, *.ucf,
-# *.qsf, *.as and *.js.
+# *.h) to filter out the source-files in the directories.
+#
+# Note that for custom extensions or not directly supported extensions you also
+# need to set EXTENSION_MAPPING for the extension otherwise the files are not
+# read by doxygen.
+#
+# If left blank the following patterns are tested:*.c, *.cc, *.cxx, *.cpp,
+# *.c++, *.java, *.ii, *.ixx, *.ipp, *.i++, *.inl, *.idl, *.ddl, *.odl, *.h,
+# *.hh, *.hxx, *.hpp, *.h++, *.cs, *.d, *.php, *.php4, *.php5, *.phtml, *.inc,
+# *.m, *.markdown, *.md, *.mm, *.dox (to be provided as doxygen C comment),
+# *.doc (to be provided as doxygen C comment), *.txt (to be provided as doxygen
+# C comment), *.py, *.pyw, *.f90, *.f95, *.f03, *.f08, *.f, *.for, *.tcl, *.vhd,
+# *.vhdl, *.ucf, *.qsf and *.ice.
 
-FILE_PATTERNS          = *.h *.cpp
+FILE_PATTERNS          = *.c \
+                         *.cc \
+                         *.cxx \
+                         *.cpp \
+                         *.h \
+                         *.hpp \
+                         *.md \
+                         *.dox \
+                         *.doc \
+                         *.txt \
 
 # The RECURSIVE tag can be used to specify whether or not subdirectories should
 # be searched for input files as well.
 # The default value is: NO.
 
-RECURSIVE              = NO
+RECURSIVE              = YES
 
 # The EXCLUDE tag can be used to specify files and/or directories that should be
 # excluded from the INPUT source files. This way you can easily exclude a
@@ -866,6 +957,10 @@ IMAGE_PATH             = $(HIP_PATH)/docs/doxygen-input/images
 # Note that the filter must not add or remove lines; it is applied before the
 # code is scanned, but not when the output code is generated. If lines are added
 # or removed, the anchors will not be placed correctly.
+#
+# Note that for custom extensions or not directly supported extensions you also
+# need to set EXTENSION_MAPPING for the extension otherwise the files are not
+# properly processed by doxygen.
 
 INPUT_FILTER           =
 
@@ -875,6 +970,10 @@ INPUT_FILTER           =
 # (like *.cpp=my_cpp_filter). See INPUT_FILTER for further information on how
 # filters are used. If the FILTER_PATTERNS tag is empty or if none of the
 # patterns match the file name, INPUT_FILTER is applied.
+#
+# Note that for custom extensions or not directly supported extensions you also
+# need to set EXTENSION_MAPPING for the extension otherwise the files are not
+# properly processed by doxygen.
 
 FILTER_PATTERNS        =
 
@@ -898,7 +997,7 @@ FILTER_SOURCE_PATTERNS =
 # (index.html). This can be useful if you have a project on for instance GitHub
 # and want to reuse the introduction page also for the doxygen output.
 
-USE_MDFILE_AS_MAINPAGE = $(HIP_PATH)/README.md
+USE_MDFILE_AS_MAINPAGE =
 
 #---------------------------------------------------------------------------
 # Configuration options related to source browsing
@@ -927,7 +1026,7 @@ INLINE_SOURCES         = NO
 STRIP_CODE_COMMENTS    = YES
 
 # If the REFERENCED_BY_RELATION tag is set to YES then for each documented
-# function all documented functions referencing it will be listed.
+# entity all documented functions referencing it will be listed.
 # The default value is: NO.
 
 REFERENCED_BY_RELATION = NO
@@ -959,12 +1058,12 @@ SOURCE_TOOLTIPS        = YES
 # If the USE_HTAGS tag is set to YES then the references to source code will
 # point to the HTML generated by the htags(1) tool instead of doxygen built-in
 # source browser. The htags tool is part of GNU's global source tagging system
-# (see http://www.gnu.org/software/global/global.html). You will need version
+# (see https://www.gnu.org/software/global/global.html). You will need version
 # 4.8.6 or higher.
 #
 # To use it do the following:
 # - Install the latest version of global
-# - Enable SOURCE_BROWSER and USE_HTAGS in the config file
+# - Enable SOURCE_BROWSER and USE_HTAGS in the configuration file
 # - Make sure the INPUT points to the root of the source tree
 # - Run doxygen as normal
 #
@@ -992,7 +1091,7 @@ VERBATIM_HEADERS       = YES
 # rich C++ code for which doxygen's built-in parser lacks the necessary type
 # information.
 # Note: The availability of this option depends on whether or not doxygen was
-# compiled with the --with-libclang option.
+# generated with the -Duse_libclang=ON option for CMake.
 # The default value is: NO.
 
 CLANG_ASSISTED_PARSING = NO
@@ -1005,6 +1104,16 @@ CLANG_ASSISTED_PARSING = NO
 
 CLANG_OPTIONS          =
 
+# If clang assisted parsing is enabled you can provide the clang parser with the
+# path to the compilation database (see:
+# http://clang.llvm.org/docs/HowToSetupToolingForLLVM.html) used when the files
+# were built. This is equivalent to specifying the "-p" option to a clang tool,
+# such as clang-check. These options will then be passed to the parser.
+# Note: The availability of this option depends on whether or not doxygen was
+# generated with the -Duse_libclang=ON option for CMake.
+
+CLANG_DATABASE_PATH    =
+
 #---------------------------------------------------------------------------
 # Configuration options related to the alphabetical class index
 #---------------------------------------------------------------------------
@@ -1123,7 +1232,7 @@ HTML_EXTRA_FILES       =
 # The HTML_COLORSTYLE_HUE tag controls the color of the HTML output. Doxygen
 # will adjust the colors in the style sheet and background images according to
 # this color. Hue is specified as an angle on a colorwheel, see
-# http://en.wikipedia.org/wiki/Hue for more information. For instance the value
+# https://en.wikipedia.org/wiki/Hue for more information. For instance the value
 # 0 represents red, 60 is yellow, 120 is green, 180 is cyan, 240 is blue, 300
 # purple, and 360 is red again.
 # Minimum value: 0, maximum value: 359, default value: 220.
@@ -1152,12 +1261,24 @@ HTML_COLORSTYLE_GAMMA  = 80
 
 # If the HTML_TIMESTAMP tag is set to YES then the footer of each generated HTML
 # page will contain the date and time when the page was generated. Setting this
-# to NO can help when comparing the output of multiple runs.
-# The default value is: YES.
+# to YES can help to show when doxygen was last run and thus if the
+# documentation is up to date.
+# The default value is: NO.
 # This tag requires that the tag GENERATE_HTML is set to YES.
 
 HTML_TIMESTAMP         = YES
 
+# If the HTML_DYNAMIC_MENUS tag is set to YES then the generated HTML
+# documentation will contain a main index with vertical navigation menus that
+# are dynamically created via JavaScript. If disabled, the navigation index will
+# consists of multiple levels of tabs that are statically embedded in every HTML
+# page. Disable this option to support browsers that do not have JavaScript,
+# like the Qt help browser.
+# The default value is: YES.
+# This tag requires that the tag GENERATE_HTML is set to YES.
+
+HTML_DYNAMIC_MENUS     = YES
+
 # If the HTML_DYNAMIC_SECTIONS tag is set to YES then the generated HTML
 # documentation will contain sections that can be hidden and shown after the
 # page has loaded.
@@ -1181,13 +1302,13 @@ HTML_INDEX_NUM_ENTRIES = 100
 
 # If the GENERATE_DOCSET tag is set to YES, additional index files will be
 # generated that can be used as input for Apple's Xcode 3 integrated development
-# environment (see: http://developer.apple.com/tools/xcode/), introduced with
-# OSX 10.5 (Leopard). To create a documentation set, doxygen will generate a
+# environment (see: https://developer.apple.com/xcode/), introduced with OSX
+# 10.5 (Leopard). To create a documentation set, doxygen will generate a
 # Makefile in the HTML output directory. Running make will produce the docset in
 # that directory and running make install will install the docset in
 # ~/Library/Developer/Shared/Documentation/DocSets so that Xcode will find it at
-# startup. See http://developer.apple.com/tools/creatingdocsetswithdoxygen.html
-# for more information.
+# startup. See https://developer.apple.com/library/archive/featuredarticles/Doxy
+# genXcode/_index.html for more information.
 # The default value is: NO.
 # This tag requires that the tag GENERATE_HTML is set to YES.
 
@@ -1226,7 +1347,7 @@ DOCSET_PUBLISHER_NAME  = Publisher
 # If the GENERATE_HTMLHELP tag is set to YES then doxygen generates three
 # additional HTML index files: index.hhp, index.hhc, and index.hhk. The
 # index.hhp is a project file that can be read by Microsoft's HTML Help Workshop
-# (see: http://www.microsoft.com/en-us/download/details.aspx?id=21138) on
+# (see: https://www.microsoft.com/en-us/download/details.aspx?id=21138) on
 # Windows.
 #
 # The HTML Help Workshop contains a compiler that can convert all HTML output
@@ -1302,7 +1423,7 @@ QCH_FILE               =
 
 # The QHP_NAMESPACE tag specifies the namespace to use when generating Qt Help
 # Project output. For more information please see Qt Help Project / Namespace
-# (see: http://qt-project.org/doc/qt-4.8/qthelpproject.html#namespace).
+# (see: https://doc.qt.io/archives/qt-4.8/qthelpproject.html#namespace).
 # The default value is: org.doxygen.Project.
 # This tag requires that the tag GENERATE_QHP is set to YES.
 
@@ -1310,7 +1431,7 @@ QHP_NAMESPACE          = org.doxygen.Project
 
 # The QHP_VIRTUAL_FOLDER tag specifies the namespace to use when generating Qt
 # Help Project output. For more information please see Qt Help Project / Virtual
-# Folders (see: http://qt-project.org/doc/qt-4.8/qthelpproject.html#virtual-
+# Folders (see: https://doc.qt.io/archives/qt-4.8/qthelpproject.html#virtual-
 # folders).
 # The default value is: doc.
 # This tag requires that the tag GENERATE_QHP is set to YES.
@@ -1319,7 +1440,7 @@ QHP_VIRTUAL_FOLDER     = doc
 
 # If the QHP_CUST_FILTER_NAME tag is set, it specifies the name of a custom
 # filter to add. For more information please see Qt Help Project / Custom
-# Filters (see: http://qt-project.org/doc/qt-4.8/qthelpproject.html#custom-
+# Filters (see: https://doc.qt.io/archives/qt-4.8/qthelpproject.html#custom-
 # filters).
 # This tag requires that the tag GENERATE_QHP is set to YES.
 
@@ -1327,7 +1448,7 @@ QHP_CUST_FILTER_NAME   =
 
 # The QHP_CUST_FILTER_ATTRS tag specifies the list of the attributes of the
 # custom filter to add. For more information please see Qt Help Project / Custom
-# Filters (see: http://qt-project.org/doc/qt-4.8/qthelpproject.html#custom-
+# Filters (see: https://doc.qt.io/archives/qt-4.8/qthelpproject.html#custom-
 # filters).
 # This tag requires that the tag GENERATE_QHP is set to YES.
 
@@ -1335,7 +1456,7 @@ QHP_CUST_FILTER_ATTRS  =
 
 # The QHP_SECT_FILTER_ATTRS tag specifies the list of the attributes this
 # project's filter section matches. Qt Help Project / Filter Attributes (see:
-# http://qt-project.org/doc/qt-4.8/qthelpproject.html#filter-attributes).
+# https://doc.qt.io/archives/qt-4.8/qthelpproject.html#filter-attributes).
 # This tag requires that the tag GENERATE_QHP is set to YES.
 
 QHP_SECT_FILTER_ATTRS  =
@@ -1428,7 +1549,7 @@ EXT_LINKS_IN_WINDOW    = NO
 
 FORMULA_FONTSIZE       = 10
 
-# Use the FORMULA_TRANPARENT tag to determine whether or not the images
+# Use the FORMULA_TRANSPARENT tag to determine whether or not the images
 # generated for formulas are transparent PNGs. Transparent PNGs are not
 # supported properly for IE 6.0, but are supported on all modern browsers.
 #
@@ -1439,8 +1560,14 @@ FORMULA_FONTSIZE       = 10
 
 FORMULA_TRANSPARENT    = YES
 
+# The FORMULA_MACROFILE can contain LaTeX \newcommand and \renewcommand commands
+# to create new LaTeX commands to be used in formulas as building blocks. See
+# the section "Including formulas" for details.
+
+FORMULA_MACROFILE      =
+
 # Enable the USE_MATHJAX option to render LaTeX formulas using MathJax (see
-# http://www.mathjax.org) which uses client side Javascript for the rendering
+# https://www.mathjax.org) which uses client side JavaScript for the rendering
 # instead of using pre-rendered bitmaps. Use this if you do not have LaTeX
 # installed or if you want to formulas look prettier in the HTML output. When
 # enabled you may also need to install MathJax separately and configure the path
@@ -1467,11 +1594,11 @@ MATHJAX_FORMAT         = HTML-CSS
 # MATHJAX_RELPATH should be ../mathjax. The default value points to the MathJax
 # Content Delivery Network so you can quickly see the result without installing
 # MathJax. However, it is strongly recommended to install a local copy of
-# MathJax from http://www.mathjax.org before deployment.
-# The default value is: http://cdn.mathjax.org/mathjax/latest.
+# MathJax from https://www.mathjax.org before deployment.
+# The default value is: https://cdnjs.cloudflare.com/ajax/libs/mathjax/2.7.5/.
 # This tag requires that the tag USE_MATHJAX is set to YES.
 
-MATHJAX_RELPATH        = http://cdn.mathjax.org/mathjax/latest
+MATHJAX_RELPATH        = https://cdnjs.cloudflare.com/ajax/libs/mathjax/2.7.5/
 
 # The MATHJAX_EXTENSIONS tag can be used to specify one or more MathJax
 # extension names that should be enabled during MathJax rendering. For example
@@ -1510,7 +1637,7 @@ MATHJAX_CODEFILE       =
 SEARCHENGINE           = YES
 
 # When the SERVER_BASED_SEARCH tag is enabled the search engine will be
-# implemented using a web server instead of a web client using Javascript. There
+# implemented using a web server instead of a web client using JavaScript. There
 # are two flavors of web server based searching depending on the EXTERNAL_SEARCH
 # setting. When disabled, doxygen will generate a PHP script for searching and
 # an index file used by the script. When EXTERNAL_SEARCH is enabled the indexing
@@ -1529,7 +1656,7 @@ SERVER_BASED_SEARCH    = NO
 #
 # Doxygen ships with an example indexer (doxyindexer) and search engine
 # (doxysearch.cgi) which are based on the open source search engine library
-# Xapian (see: http://xapian.org/).
+# Xapian (see: https://xapian.org/).
 #
 # See the section "External Indexing and Searching" for details.
 # The default value is: NO.
@@ -1542,7 +1669,7 @@ EXTERNAL_SEARCH        = NO
 #
 # Doxygen ships with an example indexer (doxyindexer) and search engine
 # (doxysearch.cgi) which are based on the open source search engine library
-# Xapian (see: http://xapian.org/). See the section "External Indexing and
+# Xapian (see: https://xapian.org/). See the section "External Indexing and
 # Searching" for details.
 # This tag requires that the tag SEARCHENGINE is set to YES.
 
@@ -1594,21 +1721,35 @@ LATEX_OUTPUT           = latex
 # The LATEX_CMD_NAME tag can be used to specify the LaTeX command name to be
 # invoked.
 #
-# Note that when enabling USE_PDFLATEX this option is only used for generating
-# bitmaps for formulas in the HTML output, but not in the Makefile that is
-# written to the output directory.
-# The default file is: latex.
+# Note that when not enabling USE_PDFLATEX the default is latex when enabling
+# USE_PDFLATEX the default is pdflatex and when in the later case latex is
+# chosen this is overwritten by pdflatex. For specific output languages the
+# default can have been set differently, this depends on the implementation of
+# the output language.
 # This tag requires that the tag GENERATE_LATEX is set to YES.
 
-LATEX_CMD_NAME         = latex
+LATEX_CMD_NAME         =
 
 # The MAKEINDEX_CMD_NAME tag can be used to specify the command name to generate
 # index for LaTeX.
+# Note: This tag is used in the Makefile / make.bat.
+# See also: LATEX_MAKEINDEX_CMD for the part in the generated output file
+# (.tex).
 # The default file is: makeindex.
 # This tag requires that the tag GENERATE_LATEX is set to YES.
 
 MAKEINDEX_CMD_NAME     = makeindex
 
+# The LATEX_MAKEINDEX_CMD tag can be used to specify the command name to
+# generate index for LaTeX. In case there is no backslash (\) as first character
+# it will be automatically added in the LaTeX code.
+# Note: This tag is used in the generated output file (.tex).
+# See also: MAKEINDEX_CMD_NAME for the part in the Makefile / make.bat.
+# The default value is: makeindex.
+# This tag requires that the tag GENERATE_LATEX is set to YES.
+
+LATEX_MAKEINDEX_CMD    = makeindex
+
 # If the COMPACT_LATEX tag is set to YES, doxygen generates more compact LaTeX
 # documents. This may be useful for small projects and may help to save some
 # trees in general.
@@ -1627,9 +1768,12 @@ COMPACT_LATEX          = NO
 PAPER_TYPE             = a4
 
 # The EXTRA_PACKAGES tag can be used to specify one or more LaTeX package names
-# that should be included in the LaTeX output. To get the times font for
-# instance you can specify
-# EXTRA_PACKAGES=times
+# that should be included in the LaTeX output. The package can be specified just
+# by its name or with the correct syntax as to be used with the LaTeX
+# \usepackage command. To get the times font for instance you can specify :
+# EXTRA_PACKAGES=times or EXTRA_PACKAGES={times}
+# To use the option intlimits with the amsmath package you can specify:
+# EXTRA_PACKAGES=[intlimits]{amsmath}
 # If left blank no extra packages will be included.
 # This tag requires that the tag GENERATE_LATEX is set to YES.
 
@@ -1671,7 +1815,7 @@ LATEX_FOOTER           =
 # list).
 # This tag requires that the tag GENERATE_LATEX is set to YES.
 
-#unsupported-LATEX_EXTRA_STYLESHEET =
+LATEX_EXTRA_STYLESHEET =
 
 # The LATEX_EXTRA_FILES tag can be used to specify one or more extra images or
 # other source files which should be copied to the LATEX_OUTPUT output
@@ -1726,12 +1870,28 @@ LATEX_SOURCE_CODE      = NO
 
 # The LATEX_BIB_STYLE tag can be used to specify the style to use for the
 # bibliography, e.g. plainnat, or ieeetr. See
-# http://en.wikipedia.org/wiki/BibTeX and \cite for more info.
+# https://en.wikipedia.org/wiki/BibTeX and \cite for more info.
 # The default value is: plain.
 # This tag requires that the tag GENERATE_LATEX is set to YES.
 
 LATEX_BIB_STYLE        = plain
 
+# If the LATEX_TIMESTAMP tag is set to YES then the footer of each generated
+# page will contain the date and time when the page was generated. Setting this
+# to NO can help when comparing the output of multiple runs.
+# The default value is: NO.
+# This tag requires that the tag GENERATE_LATEX is set to YES.
+
+LATEX_TIMESTAMP        = NO
+
+# The LATEX_EMOJI_DIRECTORY tag is used to specify the (relative or absolute)
+# path from which the emoji images will be read. If a relative path is entered,
+# it will be relative to the LATEX_OUTPUT directory. If left blank the
+# LATEX_OUTPUT directory will be used.
+# This tag requires that the tag GENERATE_LATEX is set to YES.
+
+LATEX_EMOJI_DIRECTORY  =
+
 #---------------------------------------------------------------------------
 # Configuration options related to the RTF output
 #---------------------------------------------------------------------------
@@ -1771,19 +1931,19 @@ COMPACT_RTF            = NO
 
 RTF_HYPERLINKS         = NO
 
-# Load stylesheet definitions from file. Syntax is similar to doxygen's config
-# file, i.e. a series of assignments. You only have to provide replacements,
-# missing definitions are set to their default value.
+# Load stylesheet definitions from file. Syntax is similar to doxygen's
+# configuration file, i.e. a series of assignments. You only have to provide
+# replacements, missing definitions are set to their default value.
 #
 # See also section "Doxygen usage" for information on how to generate the
 # default style sheet that doxygen normally uses.
 # This tag requires that the tag GENERATE_RTF is set to YES.
 
-#unsupported-RTF_STYLESHEET_FILE    =
+RTF_STYLESHEET_FILE    =
 
 # Set optional variables used in the generation of an RTF document. Syntax is
-# similar to doxygen's config file. A template extensions file can be generated
-# using doxygen -e rtf extensionFile.
+# similar to doxygen's configuration file. A template extensions file can be
+# generated using doxygen -e rtf extensionFile.
 # This tag requires that the tag GENERATE_RTF is set to YES.
 
 RTF_EXTENSIONS_FILE    =
@@ -1831,7 +1991,7 @@ MAN_EXTENSION          = .3
 # MAN_EXTENSION with the initial . removed.
 # This tag requires that the tag GENERATE_MAN is set to YES.
 
-#unsupported-MAN_SUBDIR             =
+MAN_SUBDIR             =
 
 # If the MAN_LINKS tag is set to YES and doxygen generates man output, then it
 # will generate one additional man file for each entity documented in the real
@@ -1869,6 +2029,13 @@ XML_OUTPUT             = xml
 
 XML_PROGRAMLISTING     = YES
 
+# If the XML_NS_MEMB_FILE_SCOPE tag is set to YES, doxygen will include
+# namespace members in file scope as well, matching the HTML output.
+# The default value is: NO.
+# This tag requires that the tag GENERATE_XML is set to YES.
+
+XML_NS_MEMB_FILE_SCOPE = NO
+
 #---------------------------------------------------------------------------
 # Configuration options related to the DOCBOOK output
 #---------------------------------------------------------------------------
@@ -1894,16 +2061,16 @@ DOCBOOK_OUTPUT         = docbook
 # The default value is: NO.
 # This tag requires that the tag GENERATE_DOCBOOK is set to YES.
 
-#unsupported-DOCBOOK_PROGRAMLISTING = NO
+DOCBOOK_PROGRAMLISTING = NO
 
 #---------------------------------------------------------------------------
 # Configuration options for the AutoGen Definitions output
 #---------------------------------------------------------------------------
 
 # If the GENERATE_AUTOGEN_DEF tag is set to YES, doxygen will generate an
-# AutoGen Definitions (see http://autogen.sf.net) file that captures the
-# structure of the code including all documentation. Note that this feature is
-# still experimental and incomplete at the moment.
+# AutoGen Definitions (see http://autogen.sourceforge.net/) file that captures
+# the structure of the code including all documentation. Note that this feature
+# is still experimental and incomplete at the moment.
 # The default value is: NO.
 
 GENERATE_AUTOGEN_DEF   = NO
@@ -2070,12 +2237,6 @@ EXTERNAL_GROUPS        = YES
 
 EXTERNAL_PAGES         = YES
 
-# The PERL_PATH should be the absolute path and name of the perl script
-# interpreter (i.e. the result of 'which perl').
-# The default file (with absolute path) is: /usr/bin/perl.
-
-PERL_PATH              = /usr/bin/perl
-
 #---------------------------------------------------------------------------
 # Configuration options related to the dot tool
 #---------------------------------------------------------------------------
@@ -2089,15 +2250,6 @@ PERL_PATH              = /usr/bin/perl
 
 CLASS_DIAGRAMS         = YES
 
-# You can define message sequence charts within doxygen comments using the \msc
-# command. Doxygen will then run the mscgen tool (see:
-# http://www.mcternan.me.uk/mscgen/)) to produce the chart and insert it in the
-# documentation. The MSCGEN_PATH tag allows you to specify the directory where
-# the mscgen tool resides. If left empty the tool is assumed to be found in the
-# default search path.
-
-MSCGEN_PATH            =
-
 # You can include diagrams made with dia in doxygen documentation. Doxygen will
 # then run dia to produce the diagram and insert it in the documentation. The
 # DIA_PATH tag allows you to specify the directory where the dia binary resides.
@@ -2116,9 +2268,9 @@ HIDE_UNDOC_RELATIONS   = YES
 # http://www.graphviz.org/), a graph visualization toolkit from AT&T and Lucent
 # Bell Labs. The other options in this section have no effect if this option is
 # set to NO
-# The default value is: NO.
+# The default value is: YES.
 
-HAVE_DOT               = NO
+HAVE_DOT               = YES
 
 # The DOT_NUM_THREADS specifies the number of dot invocations doxygen is allowed
 # to run in parallel. When set to 0 doxygen will base this on the number of
@@ -2230,7 +2382,8 @@ INCLUDED_BY_GRAPH      = YES
 #
 # Note that enabling this option will significantly increase the time of a run.
 # So in most cases it will be better to enable call graphs for selected
-# functions only using the \callgraph command.
+# functions only using the \callgraph command. Disabling a call graph can be
+# accomplished by means of the command \hidecallgraph.
 # The default value is: NO.
 # This tag requires that the tag HAVE_DOT is set to YES.
 
@@ -2241,7 +2394,8 @@ CALL_GRAPH             = NO
 #
 # Note that enabling this option will significantly increase the time of a run.
 # So in most cases it will be better to enable caller graphs for selected
-# functions only using the \callergraph command.
+# functions only using the \callergraph command. Disabling a caller graph can be
+# accomplished by means of the command \hidecallergraph.
 # The default value is: NO.
 # This tag requires that the tag HAVE_DOT is set to YES.
 
@@ -2264,11 +2418,17 @@ GRAPHICAL_HIERARCHY    = YES
 DIRECTORY_GRAPH        = YES
 
 # The DOT_IMAGE_FORMAT tag can be used to set the image format of the images
-# generated by dot.
+# generated by dot. For an explanation of the image formats see the section
+# output formats in the documentation of the dot tool (Graphviz (see:
+# http://www.graphviz.org/)).
 # Note: If you choose svg you need to set HTML_FILE_EXTENSION to xhtml in order
 # to make the SVG files visible in IE 9+ (other browsers do not have this
 # requirement).
-# Possible values are: png, jpg, gif and svg.
+# Possible values are: png, png:cairo, png:cairo:cairo, png:cairo:gd, png:gd,
+# png:gd:gd, jpg, jpg:cairo, jpg:cairo:gd, jpg:gd, jpg:gd:gd, gif, gif:cairo,
+# gif:cairo:gd, gif:gd, gif:gd:gd, svg, png:gd, png:gd:gd, png:cairo,
+# png:cairo:gd, png:cairo:cairo, png:cairo:gdiplus, png:gdiplus and
+# png:gdiplus:gdiplus.
 # The default value is: png.
 # This tag requires that the tag HAVE_DOT is set to YES.
 
@@ -2317,12 +2477,17 @@ DIAFILE_DIRS           =
 # generate a warning when it encounters a \startuml command in this case and
 # will not generate output for the diagram.
 
-#unsupported-PLANTUML_JAR_PATH      =
+PLANTUML_JAR_PATH      =
+
+# When using plantuml, the PLANTUML_CFG_FILE tag can be used to specify a
+# configuration file for plantuml.
+
+PLANTUML_CFG_FILE      =
 
 # When using plantuml, the specified paths are searched for files specified by
 # the !include statement in a plantuml block.
 
-#unsupported-PLANTUML_INCLUDE_PATH  =
+PLANTUML_INCLUDE_PATH  =
 
 # The DOT_GRAPH_MAX_NODES tag can be used to set the maximum number of nodes
 # that will be shown in the graph. If the number of nodes in a graph becomes

projects/hip/docs/doxygen-input/mainpage.txt

@@ -7,11 +7,11 @@
  * The HIP interface makes it very easy to port existing CUDA apps to run on AMD GPUs,
  * or to develop new apps that can run on either CUDA or AMD GPUs from a common source base.
  *
- *  - HIP is very thin and has little or no performance impact over coding directly in CUDA NVCC or HCC HC mode.
+ *  - HIP is very thin and has little or no performance impact over coding directly in CUDA NVCC mode.
  *  - HIP allows developers to use the "best" development environment and tools on each target platform.
  *  - HIP allows coding in a single-source C++ programming language including features such as templates, C++11 lambdas,and more.
- *  - "HIPIFY" tools automatically convert CUDA sources to HIP.
- *  - Developers can specialize for CUDA or HCC to tune for performance or handle tricky cases with #ifdef.
+ *  - HIPIFY tools automatically convert CUDA sources to HIP.
+ *  - Developers can specialize for CUDA or HIP to tune for performance or handle tricky cases with #ifdef.
 
  *  - See the @ref API.
 

projects/hip/docs/markdown/hip_porting_guide.md

@@ -496,6 +496,22 @@ int main()
 }
 ```
 
+## CU_POINTER_ATTRIBUTE_MEMORY_TYPE
+To get pointer's memory type in HIP/HIP-Clang one should use hipPointerGetAttributes API. First parameter of the API is hipPointerAttribute_t which has 'memoryType' as member variable. 'memoryType' indicates input pointer is allocated on device or host.
+
+For example:
+```
+double * ptr;
+hipMalloc(reinterpret_cast<void**>(&ptr), sizeof(double));
+hipPointerAttribute_t attr;
+hipPointerGetAttributes(&attr, ptr); /*attr.memoryType will have value as hipMemoryTypeDevice*/
+
+double* ptrHost;
+hipHostMalloc(&ptrHost, sizeof(double));
+hipPointerAttribute_t attr;
+hipPointerGetAttributes(&attr, ptrHost); /*attr.memoryType will have value as hipMemoryTypeHost*/
+```
+
 ## threadfence_system
 Threadfence_system makes all device memory writes, all writes to mapped host memory, and all writes to peer memory visible to CPU and other GPU devices.
 Some implementations can provide this behavior by flushing the GPU L2 cache.

projects/hip/include/hip/hcc_detail/hip_runtime.h

@@ -444,96 +444,111 @@ void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
 #endif
 
 #include <hip/hip_runtime_api.h>
+extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_id(uint);
+extern "C" __device__ __attribute__((const)) size_t __ockl_get_group_id(uint);
+extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_size(uint);
+extern "C" __device__ __attribute__((const)) size_t __ockl_get_num_groups(uint);
+struct __HIP_BlockIdx {
+  __device__
+  std::uint32_t operator()(std::uint32_t x) const noexcept { return __ockl_get_group_id(x); }
+};
+struct __HIP_BlockDim {
+  __device__
+  std::uint32_t operator()(std::uint32_t x) const noexcept {
+    return __ockl_get_local_size(x);
+  }
+};
+struct __HIP_GridDim {
+  __device__
+  std::uint32_t operator()(std::uint32_t x) const noexcept {
+    return __ockl_get_num_groups(x);
+  }
+};
+struct __HIP_ThreadIdx {
+  __device__
+  std::uint32_t operator()(std::uint32_t x) const noexcept {
+    return __ockl_get_local_id(x);
+  }
+};
 
-#pragma push_macro("__DEVICE__")
-#define __DEVICE__ static __device__ __forceinline__
+template <typename F>
+struct __HIP_Coordinates {
+  using R = decltype(F{}(0));
+
+  struct X { __device__ operator R() const noexcept { return F{}(0); } };
+  struct Y { __device__ operator R() const noexcept { return F{}(1); } };
+  struct Z { __device__ operator R() const noexcept { return F{}(2); } };
+
+  static constexpr X x{};
+  static constexpr Y y{};
+  static constexpr Z z{};
+#ifdef __cplusplus
+  __device__ operator dim3() const { return dim3(x, y, z); }
+#endif
+
+};
+extern "C" __device__ __attribute__((const)) size_t __ockl_get_global_size(uint);
+inline
+__device__
+std::uint32_t operator*(__HIP_Coordinates<__HIP_GridDim>::X,
+                        __HIP_Coordinates<__HIP_BlockDim>::X) noexcept {
+  return __ockl_get_global_size(0);
+}
+inline
+__device__
+std::uint32_t operator*(__HIP_Coordinates<__HIP_BlockDim>::X,
+                        __HIP_Coordinates<__HIP_GridDim>::X) noexcept {
+  return __ockl_get_global_size(0);
+}
+inline
+__device__
+std::uint32_t operator*(__HIP_Coordinates<__HIP_GridDim>::Y,
+                        __HIP_Coordinates<__HIP_BlockDim>::Y) noexcept {
+  return __ockl_get_global_size(1);
+}
+inline
+__device__
+std::uint32_t operator*(__HIP_Coordinates<__HIP_BlockDim>::Y,
+                        __HIP_Coordinates<__HIP_GridDim>::Y) noexcept {
+  return __ockl_get_global_size(1);
+}
+inline
+__device__
+std::uint32_t operator*(__HIP_Coordinates<__HIP_GridDim>::Z,
+                        __HIP_Coordinates<__HIP_BlockDim>::Z) noexcept {
+  return __ockl_get_global_size(2);
+}
+inline
+__device__
+std::uint32_t operator*(__HIP_Coordinates<__HIP_BlockDim>::Z,
+                        __HIP_Coordinates<__HIP_GridDim>::Z) noexcept {
+  return __ockl_get_global_size(2);
+}
+
+static constexpr __HIP_Coordinates<__HIP_BlockDim> blockDim{};
+static constexpr __HIP_Coordinates<__HIP_BlockIdx> blockIdx{};
+static constexpr __HIP_Coordinates<__HIP_GridDim> gridDim{};
+static constexpr __HIP_Coordinates<__HIP_ThreadIdx> threadIdx{};
 
 extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_id(uint);
-__DEVICE__ uint __hip_get_thread_idx_x() { return __ockl_get_local_id(0); }
-__DEVICE__ uint __hip_get_thread_idx_y() { return __ockl_get_local_id(1); }
-__DEVICE__ uint __hip_get_thread_idx_z() { return __ockl_get_local_id(2); }
+#define hipThreadIdx_x (__ockl_get_local_id(0))
+#define hipThreadIdx_y (__ockl_get_local_id(1))
+#define hipThreadIdx_z (__ockl_get_local_id(2))
 
 extern "C" __device__ __attribute__((const)) size_t __ockl_get_group_id(uint);
-__DEVICE__ uint __hip_get_block_idx_x() { return __ockl_get_group_id(0); }
-__DEVICE__ uint __hip_get_block_idx_y() { return __ockl_get_group_id(1); }
-__DEVICE__ uint __hip_get_block_idx_z() { return __ockl_get_group_id(2); }
+#define hipBlockIdx_x (__ockl_get_group_id(0))
+#define hipBlockIdx_y (__ockl_get_group_id(1))
+#define hipBlockIdx_z (__ockl_get_group_id(2))
 
 extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_size(uint);
-__DEVICE__ uint __hip_get_block_dim_x() { return __ockl_get_local_size(0); }
-__DEVICE__ uint __hip_get_block_dim_y() { return __ockl_get_local_size(1); }
-__DEVICE__ uint __hip_get_block_dim_z() { return __ockl_get_local_size(2); }
+#define hipBlockDim_x (__ockl_get_local_size(0))
+#define hipBlockDim_y (__ockl_get_local_size(1))
+#define hipBlockDim_z (__ockl_get_local_size(2))
 
 extern "C" __device__ __attribute__((const)) size_t __ockl_get_num_groups(uint);
-__DEVICE__ uint __hip_get_grid_dim_x() { return __ockl_get_num_groups(0); }
-__DEVICE__ uint __hip_get_grid_dim_y() { return __ockl_get_num_groups(1); }
-__DEVICE__ uint __hip_get_grid_dim_z() { return __ockl_get_num_groups(2); }
-
-#define __HIP_DEVICE_BUILTIN(DIMENSION, FUNCTION)               \
-  __declspec(property(get = __get_##DIMENSION)) uint DIMENSION; \
-  __DEVICE__ uint __get_##DIMENSION(void) {                     \
-    return FUNCTION;                                            \
-  }
-
-struct __hip_builtin_threadIdx_t {
-  __HIP_DEVICE_BUILTIN(x,__hip_get_thread_idx_x());
-  __HIP_DEVICE_BUILTIN(y,__hip_get_thread_idx_y());
-  __HIP_DEVICE_BUILTIN(z,__hip_get_thread_idx_z());
-#ifdef __cplusplus
-  __device__ operator dim3() const { return dim3(x, y, z); }
-#endif
-};
-
-struct __hip_builtin_blockIdx_t {
-  __HIP_DEVICE_BUILTIN(x,__hip_get_block_idx_x());
-  __HIP_DEVICE_BUILTIN(y,__hip_get_block_idx_y());
-  __HIP_DEVICE_BUILTIN(z,__hip_get_block_idx_z());
-#ifdef __cplusplus
-  __device__ operator dim3() const { return dim3(x, y, z); }
-#endif
-};
-
-struct __hip_builtin_blockDim_t {
-  __HIP_DEVICE_BUILTIN(x,__hip_get_block_dim_x());
-  __HIP_DEVICE_BUILTIN(y,__hip_get_block_dim_y());
-  __HIP_DEVICE_BUILTIN(z,__hip_get_block_dim_z());
-#ifdef __cplusplus
-  __device__ operator dim3() const { return dim3(x, y, z); }
-#endif
-};
-
-struct __hip_builtin_gridDim_t {
-  __HIP_DEVICE_BUILTIN(x,__hip_get_grid_dim_x());
-  __HIP_DEVICE_BUILTIN(y,__hip_get_grid_dim_y());
-  __HIP_DEVICE_BUILTIN(z,__hip_get_grid_dim_z());
-#ifdef __cplusplus
-  __device__ operator dim3() const { return dim3(x, y, z); }
-#endif
-};
-
-#undef __HIP_DEVICE_BUILTIN
-#pragma pop_macro("__DEVICE__")
-
-extern const __device__ __attribute__((weak)) __hip_builtin_threadIdx_t threadIdx;
-extern const __device__ __attribute__((weak)) __hip_builtin_blockIdx_t blockIdx;
-extern const __device__ __attribute__((weak)) __hip_builtin_blockDim_t blockDim;
-extern const __device__ __attribute__((weak)) __hip_builtin_gridDim_t gridDim;
-
-
-#define hipThreadIdx_x threadIdx.x
-#define hipThreadIdx_y threadIdx.y
-#define hipThreadIdx_z threadIdx.z
-
-#define hipBlockIdx_x blockIdx.x
-#define hipBlockIdx_y blockIdx.y
-#define hipBlockIdx_z blockIdx.z
-
-#define hipBlockDim_x blockDim.x
-#define hipBlockDim_y blockDim.y
-#define hipBlockDim_z blockDim.z
-
-#define hipGridDim_x gridDim.x
-#define hipGridDim_y gridDim.y
-#define hipGridDim_z gridDim.z
+#define hipGridDim_x (__ockl_get_num_groups(0))
+#define hipGridDim_y (__ockl_get_num_groups(1))
+#define hipGridDim_z (__ockl_get_num_groups(2))
 
 #include <hip/hcc_detail/math_functions.h>
 

projects/hip/include/hip/hcc_detail/math_functions.h

@@ -1481,6 +1481,13 @@ __DEVICE__ inline int max(int arg1, int arg2) {
   return (arg1 > arg2) ? arg1 : arg2;
 }
 
+__DEVICE__ inline int min(uint32_t arg1, int arg2) {
+  return (arg1 < arg2) ? arg1 : arg2;
+}
+__DEVICE__ inline int max(uint32_t arg1, int arg2) {
+  return (arg1 > arg2) ? arg1 : arg2;
+}
+
 __DEVICE__
 inline
 float max(float x, float y) {

projects/hip/rocclr/CMakeLists.txt

@@ -103,6 +103,7 @@ add_library(hip64 OBJECT
  hip_device_runtime.cpp
  hip_error.cpp
  hip_event.cpp
+ hip_fatbin.cpp
  hip_global.cpp
  hip_hmm.cpp
  hip_memory.cpp

projects/hip/rocclr/hip_code_object.cpp

@@ -29,15 +29,58 @@ bool CodeObject::isCompatibleCodeObject(const std::string& codeobj_target_id,
   return codeobj_target_id == short_name;
 }
 
-hipError_t CodeObject::extractCodeObjectFromFatBinary(const void* data,
-                       const std::vector<const char*>& devices,
+// This will be moved to COMGR eventually
+hipError_t CodeObject::ExtractCodeObjectFromFile(amd::Os::FileDesc fdesc, size_t fsize,
+                       const std::vector<const char*>& device_names,
                        std::vector<std::pair<const void*, size_t>>& code_objs) {
-    std::string magic((const char*)data, sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC_STR) - 1);
+
+  hipError_t hip_error = hipSuccess;
+
+  if (fdesc < 0) {
+    return hipErrorFileNotFound;
+  }
+
+  // Map the file to memory, with offset 0.
+  const void* image = nullptr;
+  if (!amd::Os::MemoryMapFileDesc(fdesc, fsize, 0, &image)) {
+    return hipErrorInvalidValue;
+  }
+
+  // retrieve code_objs{binary_image, binary_size} for devices
+  hip_error = extractCodeObjectFromFatBinary(image, device_names, code_objs);
+
+  // Unmap the file memory after extracting code object.
+  if (!amd::Os::MemoryUnmapFile(image, fsize)) {
+    return hipErrorInvalidValue;
+  }
+
+  return hip_error;
+}
+
+// This will be moved to COMGR eventually
+hipError_t CodeObject::ExtractCodeObjectFromMemory(const void* data,
+                       const std::vector<const char*>& device_names,
+                       std::vector<std::pair<const void*, size_t>>& code_objs,
+                       std::string& uri) {
+
+  // Get the URI from memory
+  if (!amd::Os::GetURIFromMemory(data, 0, uri)) {
+    return hipErrorInvalidValue;
+  }
+
+  return extractCodeObjectFromFatBinary(data, device_names, code_objs);
+}
+
+// This will be moved to COMGR eventually
+hipError_t CodeObject::extractCodeObjectFromFatBinary(const void* data,
+                       const std::vector<const char*>& device_names,
+                       std::vector<std::pair<const void*, size_t>>& code_objs) {
+  std::string magic((const char*)data, sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC_STR) - 1);
   if (magic.compare(CLANG_OFFLOAD_BUNDLER_MAGIC_STR)) {
     return hipErrorInvalidKernelFile;
   }
 
-  code_objs.resize(devices.size());
+  code_objs.resize(device_names.size());
   const auto obheader = reinterpret_cast<const __ClangOffloadBundleHeader*>(data);
   const auto* desc = &obheader->desc[0];
   unsigned num_code_objs = 0;
@@ -61,8 +104,8 @@ hipError_t CodeObject::extractCodeObjectFromFatBinary(const void* data,
         reinterpret_cast<uintptr_t>(obheader) + desc->offset);
     size_t size = desc->size;
 
-    for (size_t dev = 0; dev < devices.size(); ++dev) {
-      const char* name = devices[dev];
+    for (size_t dev = 0; dev < device_names.size(); ++dev) {
+      const char* name = device_names[dev];
 
       if (!isCompatibleCodeObject(target, name)) {
           continue;
@@ -71,7 +114,7 @@ hipError_t CodeObject::extractCodeObjectFromFatBinary(const void* data,
       num_code_objs++;
     }
   }
-  if (num_code_objs == devices.size()) {
+  if (num_code_objs == device_names.size()) {
     return hipSuccess;
   } else {
     guarantee(false && "hipErrorNoBinaryForGpu: Coudn't find binary for current devices!");
@@ -79,62 +122,31 @@ hipError_t CodeObject::extractCodeObjectFromFatBinary(const void* data,
   }
 }
 
-hipError_t CodeObject::add_program(int deviceId, hipModule_t hmod, const void* binary_ptr,
-                                   size_t binary_size) {
-  amd::Program* program = as_amd(reinterpret_cast<cl_program>(hmod));
-  amd::Context* ctx = g_devices[deviceId]->asContext();
-  if (CL_SUCCESS != program->addDeviceProgram(*ctx->devices()[0], binary_ptr,
-                                              binary_size, false)) {
-      return hipErrorNotFound;
-  }
-  return hipSuccess;
-}
-
-hipError_t CodeObject::build_module(hipModule_t hmod, const std::vector<amd::Device*>& devices) {
-  amd::Program* program = as_amd(reinterpret_cast<cl_program>(hmod));
-  program->setVarInfoCallBack(&getSvarInfo);
-  if (CL_SUCCESS != program->build(devices, nullptr, nullptr, nullptr, kOptionChangeable, kNewDevProg)) {
-    DevLogPrintfError("Build error for module: 0x%x \n", hmod);
-    return hipErrorSharedObjectInitFailed;
-  }
-  return hipSuccess;
-}
-
-DynCO::DynCO(): program_(nullptr) {}
-
 hipError_t DynCO::loadCodeObject(const char* fname, const void* image) {
 
   amd::ScopedLock lock(dclock_);
 
-  const void *mmap_ptr = nullptr;
-  size_t mmap_size = 0;
+  // Number of devices = 1 in dynamic code object
+  fb_info_ = new FatBinaryInfo(fname, image);
+  std::vector<hip::Device*> devices = { g_devices[ihipGetDevice()] };
+  IHIP_RETURN_ONFAIL(fb_info_->ExtractFatBinary(devices));
 
-  guarantee((fname || image) && "Both filename or image are nullptr");
-  if (fname != nullptr) {
-    /* We are given file name */
+  // No Lazy loading for DynCO
+  IHIP_RETURN_ONFAIL(fb_info_->BuildProgram(ihipGetDevice()));
 
-    if (!amd::Os::MemoryMapFile(fname, &mmap_ptr, &mmap_size)) {
-      return hipErrorFileNotFound;
-    }
-  } else if (image != nullptr) {
-    /*We are directly given image pointer directly */
-    mmap_ptr = image;
-  } else {
-    return hipErrorMissingConfiguration;
-  }
+  // Define Global variables
+  IHIP_RETURN_ONFAIL(populateDynGlobalVars());
 
-  return loadCodeObjectData(mmap_ptr, mmap_size);
+  // Define Global functions
+  IHIP_RETURN_ONFAIL(populateDynGlobalFuncs());
+
+  return hipSuccess;
 }
 
 //Dynamic Code Object
 DynCO::~DynCO() {
   amd::ScopedLock lock(dclock_);
 
-  if (program_ != nullptr) {
-    program_->release();
-    program_ = nullptr;
-  }
-
   for (auto& elem : vars_) {
     delete elem.second;
   }
@@ -144,6 +156,8 @@ DynCO::~DynCO() {
     delete elem.second;
   }
   functions_.clear();
+
+  delete fb_info_;
 }
 
 hipError_t DynCO::getDeviceVar(DeviceVar** dvar, std::string var_name, int device_id) {
@@ -169,52 +183,7 @@ hipError_t DynCO::getDynFunc(hipFunction_t* hfunc, std::string func_name) {
   }
 
   /* See if this could be solved */
-  return it->second->getDynFunc(hfunc, reinterpret_cast<hipModule_t>(as_cl(program_)));
-}
-
-hipError_t DynCO::loadCodeObjectData(const void* mmap_ptr, size_t mmap_size) {
-
-  amd::ScopedLock lock(dclock_);
-
-  /* initialize image it to the mmap_ptr, if this is of no_clang_offload
-     bundle then they directly pass the image */
-  const void* image = mmap_ptr;
-  std::vector<std::pair<const void*, size_t>> code_objs;
-  hipError_t hip_error = extractCodeObjectFromFatBinary(mmap_ptr,
-                            {hip::getCurrentDevice()->devices()[0]->info().name_},
-                            code_objs);
-  if (hip_error == hipSuccess) {
-    image = code_objs[0].first;
-  } else if(hip_error == hipErrorNoBinaryForGpu) {
-     return hip_error;
-  }
-
-  program_ = new amd::Program(*hip::getCurrentDevice()->asContext(),
-                              amd::Program::Language::Binary, mmap_ptr, mmap_size);
-  if (program_ == NULL) {
-    return hipErrorOutOfMemory;
-  }
-
-  program_->setVarInfoCallBack(&getSvarInfo);
-
-  if (CL_SUCCESS != program_->addDeviceProgram(*hip::getCurrentDevice()->devices()[0], image,
-                                               ElfSize(image), false)) {
-    return hipErrorInvalidKernelFile;
-  }
-
-  //This has to happen before Program has been built, other wise undef vars fail.
-  IHIP_RETURN_ONFAIL(populateDynGlobalVars());
-
-  //program->setVarInfoCallBack(&getSvarInfo);
-  if(CL_SUCCESS != program_->build(hip::getCurrentDevice()->devices(), nullptr, nullptr, nullptr,
-                                   kOptionChangeable, kNewDevProg)) {
-    return hipErrorSharedObjectInitFailed;
-  }
-
-  //This has to happen after Program has been built, other wise symbolTable_ not populated.
-  IHIP_RETURN_ONFAIL(populateDynGlobalFuncs());
-
-  return hipSuccess;
+  return it->second->getDynFunc(hfunc, module());
 }
 
 hipError_t DynCO::populateDynGlobalVars() {
@@ -223,8 +192,10 @@ hipError_t DynCO::populateDynGlobalVars() {
   std::vector<std::string> var_names;
   std::vector<std::string> undef_var_names;
 
+  //For Dynamic Modules there is only one hipFatBinaryDevInfo_
   device::Program* dev_program
-    = program_->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]);
+    = fb_info_->GetProgram(ihipGetDevice())->getDeviceProgram
+                          (*hip::getCurrentDevice()->devices()[0]);
 
   if (!dev_program->getGlobalVarFromCodeObj(&var_names)) {
     DevLogPrintfError("Could not get Global vars from Code Obj for Module: 0x%x \n", module());
@@ -252,7 +223,8 @@ hipError_t DynCO::populateDynGlobalFuncs() {
 
   std::vector<std::string> func_names;
   device::Program* dev_program
-    = program_->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]);
+    = fb_info_->GetProgram(ihipGetDevice())->getDeviceProgram(
+                           *hip::getCurrentDevice()->devices()[0]);
 
   // Get all the global func names from COMGR
   if (!dev_program->getGlobalFuncFromCodeObj(&func_names)) {
@@ -285,36 +257,21 @@ StatCO::~StatCO() {
   vars_.clear();
 }
 
-hipError_t StatCO::digestFatBinary(const void* data, FatBinaryInfoType& programs) {
+hipError_t StatCO::digestFatBinary(const void* data, FatBinaryInfo*& programs) {
   amd::ScopedLock lock(sclock_);
 
-  if (programs.size() > 0) {
+  if (programs != nullptr) {
     return hipSuccess;
   }
 
-  std::vector<std::pair<const void*, size_t>> code_objs;
-  std::vector<const char*> devices;
-  for (size_t dev = 0; dev < g_devices.size(); ++dev) {
-    devices.push_back(g_devices[dev]->devices()[0]->info().name_);
-  }
-
-  IHIP_RETURN_ONFAIL(extractCodeObjectFromFatBinary((char*)data, devices, code_objs));
-  programs.resize(g_devices.size());
-
-  for (size_t dev = 0; dev < g_devices.size(); ++dev) {
-    amd::Context* ctx = g_devices[dev]->asContext();
-    amd::Program* program = new amd::Program(*ctx);
-    if (program == nullptr) {
-      return hipErrorOutOfMemory;
-    }
-    programs.at(dev) = std::make_pair(reinterpret_cast<hipModule_t>(as_cl(program)),
-                       new FatBinaryMetaInfo(false, code_objs[dev].first, code_objs[dev].second));
-  }
+  // Create a new fat binary object and extract the fat binary for all devices.
+  programs = new FatBinaryInfo(nullptr, data);
+  IHIP_RETURN_ONFAIL(programs->ExtractFatBinary(g_devices));
 
   return hipSuccess;
 }
 
-FatBinaryInfoType* StatCO::addFatBinary(const void* data, bool initialized) {
+FatBinaryInfo** StatCO::addFatBinary(const void* data, bool initialized) {
   amd::ScopedLock lock(sclock_);
 
   if (initialized) {
@@ -324,7 +281,7 @@ FatBinaryInfoType* StatCO::addFatBinary(const void* data, bool initialized) {
   return &modules_[data];
 }
 
-hipError_t StatCO::removeFatBinary(FatBinaryInfoType* module) {
+hipError_t StatCO::removeFatBinary(FatBinaryInfo** module) {
   amd::ScopedLock lock(sclock_);
 
   auto vit = vars_.begin();
@@ -350,9 +307,7 @@ hipError_t StatCO::removeFatBinary(FatBinaryInfoType* module) {
   auto mit = modules_.begin();
   while (mit != modules_.end()) {
     if (&mit->second == module) {
-      for (size_t dev=0; dev < g_devices.size(); ++dev) {
-        delete (*module)[dev].second;
-      }
+      delete mit->second;
       mit = modules_.erase(mit);
     } else {
       ++mit;

projects/hip/rocclr/hip_code_object.hpp

@@ -21,11 +21,6 @@ class CodeObject {
 public:
   virtual ~CodeObject() {}
 
-  //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);
-
   //ClangOFFLOADBundle info
   #define CLANG_OFFLOAD_BUNDLER_MAGIC_STR "__CLANG_OFFLOAD_BUNDLE__"
   #define HIP_AMDGCN_AMDHSA_TRIPLE "hip-amdgcn-amd-amdhsa"
@@ -45,18 +40,30 @@ public:
     __ClangOffloadBundleDesc desc[1];
   };
 
+
+  // 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
+  static hipError_t ExtractCodeObjectFromFile(amd::Os::FileDesc fdesc, size_t fsize,
+                    const std::vector<const char*>& device_names,
+                    std::vector<std::pair<const void*, size_t>>& code_objs);
+
+  // Given an ptr to memory, extracts to code object for corresponding devices,
+  // returns code_objs{binary_ptr, binary_size} and uniform resource indicator
+  static hipError_t ExtractCodeObjectFromMemory(const void* data,
+                    const std::vector<const char*>& device_names,
+                    std::vector<std::pair<const void*, size_t>>& code_objs,
+                    std::string& uri);
+
+  static uint64_t ElfSize(const void* emi);
+
 protected:
+  static hipError_t extractCodeObjectFromFatBinary(const void*,
+                    const std::vector<const char*>&,
+                    std::vector<std::pair<const void*, size_t>>&);
+
   CodeObject() {}
-  //Given an ptr to image or file, extracts to code object
-  //for corresponding devices
-  hipError_t extractCodeObjectFromFatBinary(const void*,
-             const std::vector<const char*>&,
-             std::vector<std::pair<const void*, size_t>>&);
-
-  uint64_t ElfSize(const void* emi);
-
 private:
-  bool isCompatibleCodeObject(const std::string& codeobj_target_id,
+  static bool isCompatibleCodeObject(const std::string& codeobj_target_id,
                                      const char* device_name);
 
   friend const std::vector<hipModule_t>& modules();
@@ -64,30 +71,27 @@ private:
 
 //Dynamic Code Object
 class DynCO : public CodeObject {
-  amd::Monitor dclock_{"Guards Static Code object", true};
+  amd::Monitor dclock_{"Guards Dynamic Code object", true};
 
 public:
-  DynCO();
+  DynCO() {}
   virtual ~DynCO();
 
   //LoadsCodeObject and its data
   hipError_t loadCodeObject(const char* fname, const void* image=nullptr);
-  hipModule_t module() { return reinterpret_cast<hipModule_t>(as_cl(program_)); };
+  hipModule_t module() { return fb_info_->Module(ihipGetDevice()); };
 
   //Gets GlobalVar/Functions from a dynamically loaded code object
   hipError_t getDynFunc(hipFunction_t* hfunc, std::string func_name);
   hipError_t getDeviceVar(DeviceVar** dvar, std::string var_name, int deviceId);
 
 private:
-  amd::Program* program_;
+  FatBinaryInfo* fb_info_;
 
   //Maps for vars/funcs, could be keyed in with std::string name
   std::unordered_map<std::string, Function*> functions_;
   std::unordered_map<std::string, Var*> vars_;
 
-  //Load Code Object Data(Vars/UndefinedVars/Funcs)
-  hipError_t loadCodeObjectData(const void* mmap_ptr, size_t mmap_size);
-
   //Populate Global Vars/Funcs from an code object(@ module_load)
   hipError_t populateDynGlobalFuncs();
   hipError_t populateDynGlobalVars();
@@ -101,9 +105,9 @@ public:
   virtual ~StatCO();
 
   //Add/Remove/Digest Fat Binaries passed to us from "__hipRegisterFatBinary"
-  FatBinaryInfoType* addFatBinary(const void* data, bool initialized);
-  hipError_t removeFatBinary(FatBinaryInfoType* module);
-  hipError_t digestFatBinary(const void* data, FatBinaryInfoType& programs);
+  FatBinaryInfo** addFatBinary(const void* data, bool initialized);
+  hipError_t removeFatBinary(FatBinaryInfo** module);
+  hipError_t digestFatBinary(const void* data, FatBinaryInfo*& programs);
 
   //Register vars/funcs given to use from __hipRegister[Var/Func]
   hipError_t registerStatFunction(const void* hostFunction, Function* func);
@@ -120,7 +124,7 @@ public:
 private:
   friend class ::PlatformState;
   //Populated during __hipRegisterFatBinary
-  std::unordered_map<const void*, FatBinaryInfoType> modules_;
+  std::unordered_map<const void*, FatBinaryInfo*> modules_;
   //Populated during __hipRegisterFuncs
   std::unordered_map<const void*, Function*> functions_;
   //Populated during __hipRegisterVars

projects/hip/rocclr/hip_fatbin.cpp

@@ -0,0 +1,158 @@
+#include "hip_fatbin.hpp"
+
+#include "hip_code_object.hpp"
+
+namespace hip {
+
+FatBinaryDeviceInfo::~FatBinaryDeviceInfo() {
+  if (program_ != nullptr) {
+    program_->release();
+    program_ = nullptr;
+  }
+}
+
+FatBinaryInfo::FatBinaryInfo(const char* fname, const void* image)
+               : fdesc_(-1), fsize_(0), image_(image), uri_(std::string()) {
+  guarantee(fname || image);
+
+  if (fname != nullptr) {
+    fname_ = std::string(fname);
+  } else {
+    fname_ = std::string();
+  }
+
+  fatbin_dev_info_.resize(g_devices.size());
+}
+
+FatBinaryInfo::~FatBinaryInfo() {
+
+  for (auto& fbd: fatbin_dev_info_) {
+    delete fbd;
+  }
+
+  if (fdesc_ > 0) {
+    if (!amd::Os::CloseFileHandle(fdesc_)) {
+      guarantee(false && "Cannot close file");
+    }
+
+    if (!amd::Os::MemoryUnmapFile(image_, fsize_)) {
+      guarantee(false && "Cannot unmap file");
+    }
+  }
+
+  fname_ = std::string();
+  fdesc_ = -1;
+  fsize_ = 0;
+  image_ = nullptr;
+  uri_ = std::string();
+}
+
+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;
+  for (size_t dev_idx = 0; dev_idx < devices.size(); ++dev_idx) {
+    device_names.push_back(devices[dev_idx]->devices()[0]->info().name_);
+  }
+
+  // We are given file name, get the file desc and file size
+  if (fname_.size() > 0) {
+    // Get File Handle & size of the file.
+    if (!amd::Os::GetFileHandle(fname_.c_str(), &fdesc_, &fsize_)) {
+      return hipErrorFileNotFound;
+    }
+
+    // Extract the code object from file
+    hip_error = CodeObject::ExtractCodeObjectFromFile(fdesc_, fsize_,
+                device_names, code_objs);
+
+    // Map the file memory, Later: only map offset, throws error in ElfMagic now.
+    if (!amd::Os::MemoryMapFileDesc(fdesc_, fsize_, 0, &image_)) {
+      return hipErrorInvalidValue;
+    }
+
+  } else if (image_ != nullptr) {
+    // We are directly given image pointer directly, try to extract file desc & file Size
+    hip_error = CodeObject::ExtractCodeObjectFromMemory(image_,
+                device_names, code_objs, uri_);
+  } else {
+    return hipErrorMissingConfiguration;
+  }
+
+  if (hip_error == hipErrorNoBinaryForGpu) {
+    guarantee(false && "hipErrorNoBinaryForGpu: Coudn't find binary for current devices!");
+    return hip_error;
+  }
+
+  if (hip_error == hipErrorInvalidKernelFile) {
+    for (size_t dev_idx = 0; dev_idx < devices.size(); ++dev_idx) {
+      // the image type is no CLANG_OFFLOAD_BUNDLER, image for current device directly passed
+      fatbin_dev_info_[devices[dev_idx]->deviceId()]
+        = new FatBinaryDeviceInfo(image_, CodeObject::ElfSize(image_), 0);
+    }
+  } else if(hip_error == hipSuccess) {
+    for (size_t dev_idx = 0; dev_idx < devices.size(); ++dev_idx) {
+      // Calculate the offset wrt binary_image and the original image
+      size_t offset_l
+        = (reinterpret_cast<address>(const_cast<void*>(code_objs[dev_idx].first))
+            - reinterpret_cast<address>(const_cast<void*>(image_)));
+
+      fatbin_dev_info_[devices[dev_idx]->deviceId()]
+        = new FatBinaryDeviceInfo(code_objs[dev_idx].first, code_objs[dev_idx].second, offset_l);
+    }
+  }
+
+  for (size_t dev_idx = 0; dev_idx < devices.size(); ++dev_idx) {
+    fatbin_dev_info_[devices[dev_idx]->deviceId()]->program_
+       = new amd::Program(*devices[dev_idx]->asContext());
+    if (fatbin_dev_info_[devices[dev_idx]->deviceId()]->program_ == NULL) {
+      return hipErrorOutOfMemory;
+    }
+  }
+
+  return hipSuccess;
+}
+
+hipError_t FatBinaryInfo::AddDevProgram(const int device_id) {
+  // Device Id bounds Check
+  DeviceIdCheck(device_id);
+
+  FatBinaryDeviceInfo* fbd_info = fatbin_dev_info_[device_id];
+  // If fat binary was already added, skip this step and return success
+  if (fbd_info->add_dev_prog_ == false) {
+    amd::Context* ctx = g_devices[device_id]->asContext();
+    if (CL_SUCCESS != fbd_info->program_->addDeviceProgram(*ctx->devices()[0],
+                                          fbd_info->binary_image_,
+                                          fbd_info->binary_size_, false,
+                                          nullptr, nullptr, fdesc_,
+                                          fbd_info->binary_offset_, uri_)) {
+      return hipErrorInvalidKernelFile;
+    }
+    fbd_info->add_dev_prog_ = true;
+  }
+  return hipSuccess;
+}
+
+hipError_t FatBinaryInfo::BuildProgram(const int device_id) {
+
+  // Device Id Check and Add DeviceProgram if not added so far
+  DeviceIdCheck(device_id);
+  IHIP_RETURN_ONFAIL(AddDevProgram(device_id));
+
+  // If Program was already built skip this step and return success
+  FatBinaryDeviceInfo* fbd_info = fatbin_dev_info_[device_id];
+  if (fbd_info->prog_built_ == false) {
+    if(CL_SUCCESS != fbd_info->program_->build(g_devices[device_id]->devices(),
+                                               nullptr, nullptr, nullptr,
+                                               kOptionChangeable, kNewDevProg)) {
+      return hipErrorSharedObjectInitFailed;
+    }
+    fbd_info->prog_built_ = true;
+  }
+  return hipSuccess;
+}
+
+} //namespace : hip

projects/hip/rocclr/hip_fatbin.hpp

@@ -1,28 +1,86 @@
 #ifndef HIP_FAT_BINARY_HPP
 #define HIP_FAT_BINARY_HPP
 
+#include "hip/hip_runtime.h"
+#include "hip/hip_runtime_api.h"
+#include "hip_internal.hpp"
+#include "platform/program.hpp"
+
 namespace hip {
 
-class FatBinaryMetaInfo {
+//Fat Binary Per Device info
+class FatBinaryDeviceInfo {
 public:
-  FatBinaryMetaInfo(bool built, const void* binary_ptr, size_t binary_size):
-                    built_(built), binary_ptr_(binary_ptr), binary_size_(binary_size) {}
-  ~FatBinaryMetaInfo() {}
+  FatBinaryDeviceInfo (const void* binary_image, size_t binary_size, size_t binary_offset)
+                      : binary_image_(binary_image), binary_size_(binary_size),
+                        binary_offset_(binary_offset), program_(nullptr),
+                        add_dev_prog_(false), prog_built_(false) {}
 
-  //Set once the mod has been built
-  void set_built() { built_ = true; }
+  ~FatBinaryDeviceInfo();
 
-  //Accessor for private vars
-  bool built() const { return built_; }
-  const void* binary_ptr() const { return binary_ptr_; }
-  size_t binary_size() const { return binary_size_; }
 private:
-  bool built_;              //Set when mod is built. Used in Lazy Binary
-  const void* binary_ptr_;  //Binary image ptr
-  size_t binary_size_;      //Binary Size
+  const void* binary_image_; // binary image ptr
+  size_t binary_size_;       // binary image size
+  size_t binary_offset_;     // image offset from original
+
+  amd::Program* program_;    // reinterpreted as hipModule_t
+  friend class FatBinaryInfo;
+
+  //Control Variables
+  bool add_dev_prog_;
+  bool prog_built_;
 };
 
-typedef std::vector<std::pair<hipModule_t, FatBinaryMetaInfo*>> FatBinaryInfoType;
+
+// Fat Binary Info
+class FatBinaryInfo {
+public:
+  FatBinaryInfo(const char* fname, const void* image);
+  ~FatBinaryInfo();
+
+  // Loads Fat binary from file or image, unbundles COs for devices.
+  hipError_t ExtractFatBinary(const std::vector<hip::Device*>& devices);
+  hipError_t AddDevProgram(const int device_id);
+  hipError_t BuildProgram(const int device_id);
+
+
+  // Device Id bounds check
+  inline void DeviceIdCheck(const int device_id) const {
+    guarantee(device_id >= 0);
+    guarantee(static_cast<size_t>(device_id) < fatbin_dev_info_.size());
+  }
+
+  // Getter Methods
+  amd::Program* GetProgram(int device_id) {
+    DeviceIdCheck(device_id);
+    return fatbin_dev_info_[device_id]->program_;
+  }
+
+  hipModule_t Module(int device_id) const {
+    DeviceIdCheck(device_id);
+    return reinterpret_cast<hipModule_t>(as_cl(fatbin_dev_info_[device_id]->program_));
+  }
+
+  hipError_t GetModule(int device_id, hipModule_t* hmod) const {
+    DeviceIdCheck(device_id);
+    *hmod = reinterpret_cast<hipModule_t>(as_cl(fatbin_dev_info_[device_id]->program_));
+    return hipSuccess;
+  }
+
+private:
+  std::string fname_;        // File name
+  amd::Os::FileDesc fdesc_;  // File descriptor
+  size_t fsize_;             // Total file size
+
+  // Even when file is passed image will be mmapped till ~desctructor.
+  const void* image_;        // Image
+
+  // Only used for FBs where image is directly passed
+  std::string uri_;          // Uniform resource indicator
+
+  // Per Device Info, like corresponding binary ptr, size.
+  std::vector<FatBinaryDeviceInfo*> fatbin_dev_info_;
+};
 
 }; /* namespace hip */
 

projects/hip/rocclr/hip_global.cpp

@@ -81,7 +81,7 @@ DeviceFunc::~DeviceFunc() {
 }
 
 //Abstract functions
-Function::Function(std::string name, FatBinaryInfoType* modules)
+Function::Function(std::string name, FatBinaryInfo** modules)
                    : name_(name), modules_(modules) {
   dFunc_.resize(g_devices.size());
 }
@@ -106,22 +106,13 @@ hipError_t Function::getDynFunc(hipFunction_t* hfunc, hipModule_t hmod) {
 
 hipError_t Function::getStatFunc(hipFunction_t* hfunc, int deviceId) {
   guarantee(modules_ != nullptr && "Module not initialized");
-  guarantee((deviceId >= 0) && "Invalid DeviceId, less than zero");
-  guarantee((static_cast<size_t>(deviceId) < modules_->size())
-             && "Invalid DeviceId, greater than no of code objects");
 
-  hipModule_t module = (*modules_)[deviceId].first;
-  FatBinaryMetaInfo* fb_meta = (*modules_)[deviceId].second;
-
-  if (!fb_meta->built()) {
-    IHIP_RETURN_ONFAIL(CodeObject::add_program(deviceId, module, fb_meta->binary_ptr(),
-                                               fb_meta->binary_size()));
-    IHIP_RETURN_ONFAIL(CodeObject::build_module(module, g_devices[deviceId]->devices()));
-    fb_meta->set_built();
-  }
+  hipModule_t hmod = nullptr;
+  (*modules_)->BuildProgram(deviceId);
+  (*modules_)->GetModule(deviceId, &hmod);
 
   if (dFunc_[deviceId] == nullptr) {
-    dFunc_[deviceId] = new DeviceFunc(name_, (*modules_)[deviceId].first);
+    dFunc_[deviceId] = new DeviceFunc(name_, hmod);
   }
   *hfunc = dFunc_[deviceId]->asHipFunction();
 
@@ -130,22 +121,13 @@ hipError_t Function::getStatFunc(hipFunction_t* hfunc, int deviceId) {
 
 hipError_t Function::getStatFuncAttr(hipFuncAttributes* func_attr, int deviceId) {
   guarantee((modules_ != nullptr) && "Module not initialized");
-  guarantee((deviceId >= 0) && "Invalid DeviceId, less than zero");
-  guarantee((static_cast<size_t>(deviceId) < modules_->size())
-            && "Invalid DeviceId, greater than no of code objects");
 
-  hipModule_t module = (*modules_)[deviceId].first;
-  FatBinaryMetaInfo* fb_meta = (*modules_)[deviceId].second;
-
-  if (!fb_meta->built()) {
-    IHIP_RETURN_ONFAIL(CodeObject::add_program(deviceId, module, fb_meta->binary_ptr(),
-                                               fb_meta->binary_size()));
-    IHIP_RETURN_ONFAIL(CodeObject::build_module(module, g_devices[deviceId]->devices()));
-    fb_meta->set_built();
-  }
+  hipModule_t hmod = nullptr;
+  (*modules_)->BuildProgram(deviceId);
+  (*modules_)->GetModule(deviceId, &hmod);
 
   if (dFunc_[deviceId] == nullptr) {
-    dFunc_[deviceId] = new DeviceFunc(name_, (*modules_)[deviceId].first);
+    dFunc_[deviceId] = new DeviceFunc(name_, hmod);
   }
 
   const std::vector<amd::Device*>& devices = amd::Device::getDevices(CL_DEVICE_TYPE_GPU, false);
@@ -171,7 +153,7 @@ hipError_t Function::getStatFuncAttr(hipFuncAttributes* func_attr, int deviceId)
 
 //Abstract Vars
 Var::Var(std::string name, DeviceVarKind dVarKind, size_t size, int type, int norm,
-         FatBinaryInfoType* modules) : name_(name), dVarKind_(dVarKind), size_(size),
+         FatBinaryInfo** modules) : name_(name), dVarKind_(dVarKind), size_(size),
          type_(type), norm_(norm), modules_(modules) {
   dVar_.resize(g_devices.size());
 }
@@ -203,18 +185,12 @@ hipError_t Var::getStatDeviceVar(DeviceVar** dvar, int deviceId) {
   guarantee((static_cast<size_t>(deviceId) < g_devices.size())
             && "Invalid DeviceId, greater than no of code objects");
 
-  hipModule_t module = (*modules_)[deviceId].first;
-  FatBinaryMetaInfo* fb_meta = (*modules_)[deviceId].second;
-
-  if (!fb_meta->built()) {
-    IHIP_RETURN_ONFAIL(CodeObject::add_program(deviceId, module, fb_meta->binary_ptr(),
-                                               fb_meta->binary_size()));
-    IHIP_RETURN_ONFAIL(CodeObject::build_module(module, g_devices[deviceId]->devices()));
-    fb_meta->set_built();
-  }
+  hipModule_t hmod = nullptr;
+  (*modules_)->BuildProgram(deviceId);
+  (*modules_)->GetModule(deviceId, &hmod);
 
   if (dVar_[deviceId] == nullptr) {
-    dVar_[deviceId] = new DeviceVar(name_, (*modules_)[deviceId].first);
+    dVar_[deviceId] = new DeviceVar(name_, hmod);
   }
 
   *dvar = dVar_[deviceId];

projects/hip/rocclr/hip_global.hpp

@@ -57,7 +57,7 @@ private:
 //Abstract Structures
 class Function {
 public:
-  Function(std::string name, FatBinaryInfoType* modules=nullptr);
+  Function(std::string name, FatBinaryInfo** modules=nullptr);
   ~Function();
 
   //Return DeviceFunc for this this dynamically loaded module
@@ -67,12 +67,12 @@ public:
   hipError_t getStatFunc(hipFunction_t *hfunc, int deviceId);
   hipError_t getStatFuncAttr(hipFuncAttributes* func_attr, int deviceId);
   void resize_dFunc(size_t size) { dFunc_.resize(size); }
-  FatBinaryInfoType* moduleInfo() { return modules_; };
+  FatBinaryInfo** moduleInfo() { return modules_; };
 
 private:
   std::vector<DeviceFunc*> dFunc_;  //DeviceFuncObj per Device
   std::string name_;                //name of the func(not unique identifier)
-  FatBinaryInfoType* modules_;      // static module where it is referenced
+  FatBinaryInfo** modules_;      // static module where it is referenced
 };
 
 class Var {
@@ -85,7 +85,7 @@ public:
   };
 
   Var(std::string name, DeviceVarKind dVarKind, size_t size, int type, int norm,
-      FatBinaryInfoType* modules = nullptr);
+      FatBinaryInfo** modules = nullptr);
   ~Var();
 
   //Return DeviceVar for this dynamically loaded module
@@ -97,10 +97,10 @@ public:
 
   //Accessor for device_ptrs.
   std::string name() const { return name_; }
-  hipModule_t module(int deviceId) const { return (*modules_)[deviceId].first; }
+  hipModule_t module(int deviceId) const { return nullptr; }
   hipDeviceptr_t device_ptr(int deviceId) const { return dVar_[deviceId]->device_ptr(); }
   size_t device_size(int deviceId) const { return dVar_[deviceId]->size(); }
-  FatBinaryInfoType* moduleInfo() { return modules_; };
+  FatBinaryInfo** moduleInfo() { return modules_; };
 
 private:
   std::vector<DeviceVar*> dVar_;   // DeviceVarObj per Device
@@ -109,7 +109,7 @@ private:
   size_t size_;                    // Size of the variable
   int type_;                       // Type(Textures/Surfaces only)
   int norm_;                       // Type(Textures/Surfaces only)
-  FatBinaryInfoType* modules_;     // static module where it is referenced
+  FatBinaryInfo** modules_;     // static module where it is referenced
 };
 
 }; //namespace: hip

projects/hip/rocclr/hip_hcc.def.in

@@ -266,3 +266,5 @@ hipMemcpy2DFromArrayAsync
 hipMemcpyAtoH
 hipMemcpyHtoA
 hipMemcpyParam2DAsync
+__gnu_h2f_ieee
+__gnu_f2h_ieee

projects/hip/rocclr/hip_hcc.map.in

@@ -178,6 +178,8 @@ global:
     __hipRegisterSurface;
     __hipRegisterTexture;
     __hipUnregisterFatBinary;
+    __gnu_h2f_ieee;
+    __gnu_f2h_ieee;
     hipConfigureCall;
     hipSetupArgument;
     hipLaunchByPtr;

projects/hip/rocclr/hip_platform.cpp

@@ -68,7 +68,7 @@ static bool isCompatibleCodeObject(const std::string& codeobj_target_id,
   return codeobj_target_id == short_name;
 }
 
-extern "C" hip::FatBinaryInfoType* __hipRegisterFatBinary(const void* data)
+extern "C" hip::FatBinaryInfo** __hipRegisterFatBinary(const void* data)
 {
   const __CudaFatBinaryWrapper* fbwrapper = reinterpret_cast<const __CudaFatBinaryWrapper*>(data);
   if (fbwrapper->magic != __hipFatMAGIC2 || fbwrapper->version != 1) {
@@ -102,7 +102,7 @@ bool CL_CALLBACK getSvarInfo(cl_program program, std::string var_name, void** va
 }
 
 extern "C" void __hipRegisterFunction(
-  hip::FatBinaryInfoType* modules,
+  hip::FatBinaryInfo** modules,
   const void*  hostFunction,
   char*        deviceFunction,
   const char*  deviceName,
@@ -137,7 +137,7 @@ extern "C" void __hipRegisterFunction(
 // track of the value of the device side global variable between kernel
 // executions.
 extern "C" void __hipRegisterVar(
-  hip::FatBinaryInfoType* modules,   // The device modules containing code object
+  hip::FatBinaryInfo** modules,   // The device modules containing code object
   void*       var,       // The shadow variable in host code
   char*       hostVar,   // Variable name in host code
   char*       deviceVar, // Variable name in device code
@@ -150,7 +150,7 @@ extern "C" void __hipRegisterVar(
   PlatformState::instance().registerStatGlobalVar(var, var_ptr);
 }
 
-extern "C" void __hipRegisterSurface(hip::FatBinaryInfoType* modules,      // The device modules containing code object
+extern "C" void __hipRegisterSurface(hip::FatBinaryInfo** modules,      // The device modules containing code object
                                      void* var,        // The shadow variable in host code
                                      char* hostVar,    // Variable name in host code
                                      char* deviceVar,  // Variable name in device code
@@ -159,7 +159,7 @@ extern "C" void __hipRegisterSurface(hip::FatBinaryInfoType* modules,      // Th
   PlatformState::instance().registerStatGlobalVar(var, var_ptr);
 }
 
-extern "C" void __hipRegisterTexture(hip::FatBinaryInfoType* modules,      // The device modules containing code object
+extern "C" void __hipRegisterTexture(hip::FatBinaryInfo** modules,      // The device modules containing code object
                                      void* var,        // The shadow variable in host code
                                      char* hostVar,    // Variable name in host code
                                      char* deviceVar,  // Variable name in device code
@@ -168,7 +168,7 @@ extern "C" void __hipRegisterTexture(hip::FatBinaryInfoType* modules,      // Th
   PlatformState::instance().registerStatGlobalVar(var, var_ptr);
 }
 
-extern "C" void __hipUnregisterFatBinary(hip::FatBinaryInfoType* modules)
+extern "C" void __hipUnregisterFatBinary(hip::FatBinaryInfo** modules)
 {
   HIP_INIT();
 
@@ -644,23 +644,69 @@ hipError_t ihipLaunchKernel(const void* hostFunction,
                                     flags));
 }
 
+// conversion routines between float and half precision
+
+static inline std::uint32_t f32_as_u32(float f) { union { float f; std::uint32_t u; } v; v.f = f; return v.u; }
+
+static inline float u32_as_f32(std::uint32_t u) { union { float f; std::uint32_t u; } v; v.u = u; return v.f; }
+
+static inline int clamp_int(int i, int l, int h) { return std::min(std::max(i, l), h); }
+
+
+// half float, the f16 is in the low 16 bits of the input argument
+
+static inline float __convert_half_to_float(std::uint32_t a) noexcept {
+
+  std::uint32_t u = ((a << 13) + 0x70000000U) & 0x8fffe000U;
+
+  std::uint32_t v = f32_as_u32(u32_as_f32(u) * u32_as_f32(0x77800000U)/*0x1.0p+112f*/) + 0x38000000U;
+
+  u = (a & 0x7fff) != 0 ? v : u;
+
+  return u32_as_f32(u) * u32_as_f32(0x07800000U)/*0x1.0p-112f*/;
+
+}
+
+// float half with nearest even rounding
+// The lower 16 bits of the result is the bit pattern for the f16
+static inline std::uint32_t __convert_float_to_half(float a) noexcept {
+  std::uint32_t u = f32_as_u32(a);
+  int e = static_cast<int>((u >> 23) & 0xff) - 127 + 15;
+  std::uint32_t m = ((u >> 11) & 0xffe) | ((u & 0xfff) != 0);
+  std::uint32_t i = 0x7c00 | (m != 0 ? 0x0200 : 0);
+  std::uint32_t n = ((std::uint32_t)e << 12) | m;
+  std::uint32_t s = (u >> 16) & 0x8000;
+  int b = clamp_int(1-e, 0, 13);
+  std::uint32_t d = (0x1000 | m) >> b;
+  d |= (d << b) != (0x1000 | m);
+  std::uint32_t v = e < 1 ? d : n;
+  v = (v >> 2) + (((v & 0x7) == 3) | ((v & 0x7) > 5));
+  v = e > 30 ? 0x7c00 : v;
+  v = e == 143 ? i : v;
+  return s | v;
+}
+
+extern "C" __attribute__((weak)) float  __gnu_h2f_ieee(unsigned short h){
+  return __convert_half_to_float((std::uint32_t) h);
+}
+
+extern "C" __attribute__((weak)) unsigned short  __gnu_f2h_ieee(float f){
+  return (unsigned short)__convert_float_to_half(f);
+}
+
 void PlatformState::init()
 {
   amd::ScopedLock lock(lock_);
-
   if(initialized_ || g_devices.empty()) {
     return;
   }
   initialized_ = true;
-
   for (auto& it : statCO_.modules_) {
     digestFatBinary(it.first, it.second);
   }
-
   for (auto &it : statCO_.vars_) {
     it.second->resize_dVar(g_devices.size());
   }
-
   for (auto &it : statCO_.functions_) {
     it.second->resize_dFunc(g_devices.size());
   }
@@ -789,15 +835,15 @@ hipError_t PlatformState::getDynTexRef(const char* hostVar, hipModule_t hmod, te
   return hipSuccess;
 }
 
-hipError_t PlatformState::digestFatBinary(const void* data, hip::FatBinaryInfoType& programs) {
+hipError_t PlatformState::digestFatBinary(const void* data, hip::FatBinaryInfo*& programs) {
  return statCO_.digestFatBinary(data, programs);
 }
 
-hip::FatBinaryInfoType* PlatformState::addFatBinary(const void* data) {
+hip::FatBinaryInfo** PlatformState::addFatBinary(const void* data) {
   return statCO_.addFatBinary(data, initialized_);
 }
 
-hipError_t PlatformState::removeFatBinary(hip::FatBinaryInfoType* module) {
+hipError_t PlatformState::removeFatBinary(hip::FatBinaryInfo** module) {
   return statCO_.removeFatBinary(module);
 }
 

projects/hip/rocclr/hip_platform.hpp

@@ -66,9 +66,9 @@ public:
   }
 
   //Static Code Objects functions
-  hip::FatBinaryInfoType* addFatBinary(const void* data);
-  hipError_t removeFatBinary(hip::FatBinaryInfoType* module);
-  hipError_t digestFatBinary(const void* data, hip::FatBinaryInfoType& programs);
+  hip::FatBinaryInfo** addFatBinary(const void* data);
+  hipError_t removeFatBinary(hip::FatBinaryInfo** module);
+  hipError_t digestFatBinary(const void* data, hip::FatBinaryInfo*& programs);
 
   hipError_t registerStatFunction(const void* hostFunction, hip::Function* func);
   hipError_t registerStatGlobalVar(const void* hostVar, hip::Var* var);

projects/hip/tests/src/cg/hipCGGridGroupType.cpp

@@ -137,6 +137,12 @@ int main()
   ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess);
   int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock;
 
+  if (!deviceProperties.cooperativeLaunch) {
+    std::cout << "info: Device doesn't support cooperative launch! skipping the test!\n";
+    passed();
+    return 0;
+  }
+
   // Test block sizes which are powers of 2
   int i = 0;
   while (true) {

projects/hip/tests/src/cg/hipCGGridGroupTypeViaBaseType.cpp

@@ -137,6 +137,12 @@ int main()
   ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess);
   int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock;
 
+  if (!deviceProperties.cooperativeLaunch) {
+    std::cout << "info: Device doesn't support cooperative launch! skipping the test!\n";
+    passed();
+    return 0;
+  }
+
   // Test block sizes which are powers of 2
   int i = 0;
   while (true) {

projects/hip/tests/src/cg/hipCGGridGroupTypeViaPublicApi.cpp

@@ -137,6 +137,12 @@ int main()
   ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess);
   int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock;
 
+  if (!deviceProperties.cooperativeLaunch) {
+    std::cout << "info: Device doesn't support cooperative launch! skipping the test!\n";
+    passed();
+    return 0;
+  }
+
   // Test block sizes which are powers of 2
   int i = 0;
   while (true) {

projects/hip/tests/src/runtimeApi/p2p/hipDeviceGetP2PAttribute.cpp

@@ -0,0 +1,53 @@
+/*
+Copyright (c) 2015-2017 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.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * TEST: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+void runTest(int srcDevice, int dstDevice) {
+  std::cout<<"Between Device "<<srcDevice<<" and Device "<<dstDevice<<std::endl;
+  for (int p2p_attr_idx = hipDevP2PAttrPerformanceRank;
+        p2p_attr_idx <= hipDevP2PAttrHipArrayAccessSupported; ++p2p_attr_idx) {
+    int value = -1;
+    HIPCHECK(hipDeviceGetP2PAttribute(&value, static_cast<hipDeviceP2PAttr>(p2p_attr_idx),
+                                      srcDevice, dstDevice));
+    std::cout<<"Attr["<<p2p_attr_idx<<"] is "<<value<<std::endl;
+  }
+}
+
+int main() {
+
+  int count = -1;
+  HIPCHECK(hipGetDeviceCount(&count))
+
+  if (count >= 2){
+    for (int dev_idx = 0; dev_idx < (count-1); ++dev_idx) {
+      runTest(dev_idx, 1 + dev_idx);
+    }
+  } else {
+    std::cout<<"Not enough GPUs to run the single GPU tests"<<std::endl;
+  }
+
+  passed();
+}