Add 'projects/hipother/' from commit 'e9d2abe72168de35233ed77167f7920509c05a8a'

git-subtree-dir: projects/hipother
git-subtree-mainline: d76041b87b
git-subtree-split: e9d2abe721

projects/hipother/.azuredevops/rocm-ci.yml

@@ -0,0 +1,50 @@
+resources:
+  repositories:
+  - repository: pipelines_repo
+    type: github
+    endpoint: ROCm
+    name: ROCm/ROCm
+  - repository: matching_repo
+    type: github
+    endpoint: ROCm
+    name: ROCm/clr
+    ref: $(Build.SourceBranch)
+  - repository: hipother_repo
+    type: github
+    endpoint: ROCm
+    name: ROCm/HIP # leverage HIP job that builds both AMD and NV backends
+    ref: $(Build.SourceBranch)
+
+variables:
+- group: common
+- template: /.azuredevops/variables-global.yml@pipelines_repo
+
+trigger:
+  batch: true
+  branches:
+    include:
+    - amd-staging
+    - amd-mainline
+  paths:
+    exclude:
+    - '.github'
+    - CODEOWNERS
+    - LICENSE.txt
+    - '**/*.md'
+
+pr:
+  autoCancel: true
+  branches:
+    include:
+    - amd-staging
+    - amd-mainline
+  paths:
+    exclude:
+    - '.github'
+    - CODEOWNERS
+    - LICENSE.txt
+    - '**/*.md'
+  drafts: false
+
+jobs:
+  - template: ${{ variables.CI_COMPONENT_PATH }}/HIP.yml@pipelines_repo

projects/hipother/.github/hooks/clang-format-check.sh

@@ -0,0 +1,56 @@
+#!/usr/bin/env bash
+
+set -euo pipefail
+
+RANGE=""
+
+while [[ $# -gt 0 ]]; do
+  echo $1
+  echo $2
+  case "$1" in
+  --range)
+    RANGE="$2"
+    shift 2
+    ;;
+  *)
+    echo "Unknown arg $1" >&2
+    exit 64
+    ;;
+  esac
+done
+
+regex='\.(c|cc|cpp|cxx|h|hh|hpp|hxx)$'
+
+if [[ -n $RANGE ]]; then
+  files=$(git diff --name-only "$RANGE" | grep -E "$regex" || true)
+else
+  files=$(git diff --cached --name-only --diff-filter=ACMR | grep -E "$regex" || true)
+fi
+echo "Checking $files"
+[[ -z $files ]] && exit 0
+
+clang_bin="${CLANG_FORMAT:-clang-format}"
+if ! command -v "$clang_bin" >/dev/null 2>&1; then
+  if [[ -x "/c/Program Files/LLVM/bin/clang-format.exe" ]]; then
+    clang_bin="/c/Program Files/LLVM/bin/clang-format.exe"
+  fi
+fi
+
+clang_format_diff="${CLANG_FORMAT_DIFF:-clang-format-diff}"
+if ! command -v "$clang_format_diff" >/dev/null 2>&1; then
+  if [[ -x "/c/Program Files/LLVM/share/clang/clang-format-diff.py" ]]; then
+    clang_format_diff="/c/Program Files/LLVM/share/clang/clang-format-diff.py"
+  fi
+fi
+
+for file in $files; do
+  echo "Checking lines of $file"
+
+  if [[ -n $RANGE ]]; then
+    diff_output=$(git diff -U0 "$RANGE" -- "$file")
+  else
+    diff_output=$(git diff -U0 --cached -- "$file")
+  fi
+
+  echo "$diff_output" | "$clang_format_diff" -style=file -fallback-style=none -p1
+done

projects/hipother/.github/hooks/pre-commit

@@ -0,0 +1,2 @@
+#!/usr/bin/env bash
+exec "$(git rev-parse --show-toplevel)/.github/hooks/clang-format-check.sh"

projects/hipother/.github/palamida.yml

@@ -0,0 +1,5 @@
+disabled: false
+scmId: gh-emu-rocm
+branchesToScan:
+  - amd-staging
+  - amd-mainline

projects/hipother/.github/pull_request_template.md

@@ -0,0 +1,36 @@
+## Associated JIRA ticket number/Github issue number
+<!-- For example: "Closes #1234" or "Fixes SWDEV-123456" -->
+
+## What type of PR is this? (check all applicable)
+
+- [ ] Refactor
+- [ ] Feature
+- [ ] Bug Fix
+- [ ] Optimization
+- [ ] Documentation Update
+- [ ] Continuous Integration
+
+## What were the changes?
+
+<!-- Please give a short summary of the change. -->
+
+## Why are these changes needed?
+
+<!-- Please explain the motivation behind the change and why this solves the given problem. -->
+
+## Updated CHANGELOG?
+
+<!-- Needed for Release updates for a ROCm release. -->
+
+- [ ] Yes
+- [ ] No, Does not apply to this PR.
+
+## Added/Updated documentation?
+
+- [ ] Yes
+- [ ] No, Does not apply to this PR.
+
+## Additional Checks
+
+- [ ] I have added tests relevant to the introduced functionality, and the unit tests are passing locally.
+- [ ] Any dependent changes have been merged.

projects/hipother/.github/scripts/validate_pr_description.py

@@ -0,0 +1,76 @@
+import os, re, sys
+from typing import List, Optional
+
+
+def is_checkbox(line: str) -> bool:
+    return bool(re.match(r"^\s*-\s*\[[ xX]\]\s*.+", line))
+
+
+def is_checked(line: str) -> bool:
+    return bool(re.match(r"^\s*-\s*\[[xX]\]\s*.+", line))
+
+
+def is_comment(line: str) -> bool:
+    return bool(re.match(r"^\s*<!--.*-->\s*$", line))
+
+
+def text_clean(lines: List[str]) -> str:
+    text = [line for line in lines if not is_comment(line)]
+    return "".join("".join(text).strip().split())
+
+
+def validate_section(section_name: str, lines: List[str]) -> Optional[str]:
+    has_checkboxes = any(is_checkbox(line) for line in lines)
+    if has_checkboxes:
+        if not any(is_checked(line) for line in lines):
+            return f"Section {section_name} is a checklist without selections"
+        return None
+    if not text_clean(lines):
+        return f"Section {section_name} is empty text section"
+    return None
+
+
+def check_description(description: str) -> List[str]:
+    if not description:
+        # pull_request_template is not merged yet, so treat as valid for now
+        return []
+        # return ["PR description is empty"]
+
+    sections = []
+    current_section = None
+    current_lines = []
+    errors = []
+
+    for line in description.splitlines():
+        header_match = re.match(r"^\s*##\s*(.+?)\s*$", line)
+        if header_match:
+            if current_section:
+                sections.append((current_section, current_lines))
+            current_section = header_match.group(1)
+            current_lines = []
+        elif current_section:
+            current_lines.append(line)
+
+    if current_section:
+        sections.append((current_section, current_lines))
+
+    if not sections:
+        return ["No sections available, template is empty"]
+
+    for section_name, section_lines in sections:
+        error = validate_section(section_name, section_lines)
+        if error:
+            errors.append(error)
+
+    return errors
+
+
+if __name__ == "__main__":
+    pr_description = os.getenv("PR_DESCRIPTION", "")
+
+    errors = check_description(pr_description)
+    if not errors:
+        print("All good")
+        exit(0)
+    print("\n".join(errors))
+    exit(1)

projects/hipother/.github/workflows/clang-format.yml

@@ -0,0 +1,22 @@
+name: Clang format check  
+on:
+  pull_request:
+    types: [synchronize, opened]
+
+jobs:
+  format:
+    runs-on: AMD-ROCm-Internal-dev1
+    steps:
+      - uses: actions/checkout@v4
+        with:
+          fetch-depth: 0
+
+      - name: Install clang-format
+        run: |
+          sudo apt update && sudo apt install -y clang-format
+
+      - name: Run clang-format-check
+        id: clang-format
+        run: |
+          chmod +x .github/hooks/clang-format-check.sh
+          ./.github/hooks/clang-format-check.sh --range "${{ github.event.pull_request.base.sha }}..${{ github.event.pull_request.head.sha }}"

projects/hipother/.github/workflows/keyword-check.yml

@@ -0,0 +1,73 @@
+name: Keywords checker
+
+on:
+  pull_request:
+    types: [opened, synchronize, reopened, edited]
+    branches:
+      - amd-staging
+  workflow_dispatch:
+
+jobs:
+  check-keywords:
+    runs-on: AMD-ROCm-Internal-dev1
+    env:
+      KEYWORDS: ${{ vars.KEYWORDS }}
+
+    steps:
+      - name: Checkout code
+        uses: actions/checkout@v3
+        with:
+          fetch-depth: 0
+
+      - name: Check keywords
+        run: |
+          set -e
+
+          if [ -z "$KEYWORDS" ]; then
+            echo "No keywords set. Skipping check"
+            exit 0
+          fi
+
+          IFS=',' read -ra KEYWORDS_ARRAY <<< "$KEYWORDS"
+          echo "Checking against list of keywords: ${KEYWORDS_ARRAY[*]}"
+
+          MATCHED=0
+          BASE_BRANCH=${{github.event.pull_request.base.ref}}
+          HEAD_BRANCH=${{github.event.pull_request.head.ref}}
+          PR_TITLE="${{ github.event.pull_request.title }}"
+
+          for file in $(git diff --name-only origin/$BASE_BRANCH..origin/$HEAD_BRANCH); do
+            if [ -f "$file" ]; then
+              for keyword in "${KEYWORDS_ARRAY[*]}"; do
+                grep -in -E "${keyword}" "$file" | while IFS= read -r line; do
+                  echo "Matched in '$file': $line"
+                  MATCHED=1
+                done
+              done
+            fi
+          done
+
+          for commit in $(git log --format=%H origin/$BASE_BRANCH..origin/$HEAD_BRANCH); do
+            msg=$(git log -1 --format=%B "$commit")
+            for keyword in "${KEYWORDS_ARRAY[*]}"; do
+              if echo "$msg" | grep -i -q "$keyword"; then
+                echo "Match in commit $commit: $msg"
+                MATCHED=1
+              fi
+            done
+          done
+
+          for keyword in "${KEYWORDS_ARRAY[*]}"; do
+            if echo "$PR_TITLE" | grep -i -q "$keyword"; then
+              echo "Match in PR title"
+              MATCHED=1
+            fi
+          done
+          
+          if [ "$MATCHED" -eq 1 ]; then
+            echo "Keywords found, please see diagnostics higher"
+            exit 1
+          else
+            echo "No keywords found"
+            exit 0
+          fi

projects/hipother/.github/workflows/kws-caller.yml

@@ -0,0 +1,15 @@
+name: Rocm Validation Suite KWS
+on:
+  push:
+    branches: [amd-staging, amd-mainline]
+  pull_request:
+    types: [opened, synchronize, reopened]
+  workflow_dispatch:
+jobs:
+  kws:
+    if: ${{ github.event_name == 'pull_request' }}
+    uses: AMD-ROCm-Internal/rocm_ci_infra/.github/workflows/kws.yml@mainline
+    secrets: inherit
+    with:
+      pr_number: ${{github.event.pull_request.number}}
+      base_branch: ${{github.base_ref}}

projects/hipother/.github/workflows/pr-title-validate.yml

@@ -0,0 +1,46 @@
+name: Validate PR Title
+
+on:
+  pull_request:
+    types: [opened, edited, synchronize, reopened]
+
+jobs:
+  validate-pr-title:
+    runs-on: ubuntu-latest
+    steps:
+      - name: Check PR Title
+        id: check-pr-title
+        run: |
+          PR_TITLE="${{ github.event.pull_request.title }}"
+
+          if [[ ! "$PR_TITLE" =~ ^SWDEV-[0-9]+ ]]; then
+            echo "::error::PR title must start with a Jira ticket ID, SWDEV-<num>"
+            exit 1
+          else
+            echo "PR title is valid"
+          fi
+
+  validate-commit-messages:
+    runs-on: AMD-ROCm-Internal-dev1
+    steps:
+      - name: Checkout code
+        uses: actions/checkout@v3
+        with:
+            fetch-depth: 0
+
+      - name: Check all commit messages
+        id: validate-commit-messags
+        run: |
+          COMMITS=$(git log --format="%H %s" origin/${{ github.event.pull_request.base.ref }}..origin/${{ github.event.pull_request.head.ref }})
+          echo "$COMMITS"
+          echo "$COMMITS" | while read -r hash message; do
+            echo -e "$hash $message\n "
+            if [[ "$message" =~ ^SWDEV-[0-9]+ ]]; then
+              echo "Valid JIRA ticket format"
+            elif [[ "$message" =~ ^Merge\ branch ]]; then
+              echo "Merge commits are allowed"
+            else
+              echo "::error:: $hash commit should start with Jira ticket ID, SWDEV-<num> or be a merge commit"
+              exit 1
+            fi
+          done

projects/hipother/.github/workflows/rocm-ci-caller.yml

@@ -0,0 +1,25 @@
+name: ROCm CI Caller  
+on:  
+  pull_request:  
+    branches: [amd-staging, release/rocm-rel-*, amd-mainline]  
+    types: [opened, reopened, synchronize]  
+  push:  
+    branches: [amd-mainline]  
+  workflow_dispatch:  
+  issue_comment:  
+    types: [created]  
+  
+jobs:  
+  call-workflow:  
+    if: github.event_name != 'issue_comment' ||(github.event_name == 'issue_comment' && github.event.issue.pull_request && (startsWith(github.event.comment.body, '!verify') || startsWith(github.event.comment.body, '!linux-hip-psdb') || startsWith(github.event.comment.body, '!verify release') || startsWith(github.event.comment.body, '!verify retest')))  
+    uses: AMD-ROCm-Internal/rocm_ci_infra/.github/workflows/rocm_ci.yml@mainline
+    secrets: inherit  
+    with:  
+      input_sha: ${{github.event_name == 'pull_request' && github.event.pull_request.head.sha || (github.event_name == 'push' && github.sha) || (github.event_name == 'issue_comment' && github.event.issue.pull_request.head.sha) || github.sha}}  
+      input_pr_num: ${{github.event_name == 'pull_request' && github.event.pull_request.number || (github.event_name == 'issue_comment' && github.event.issue.number) || 0}}  
+      input_pr_url: ${{github.event_name == 'pull_request' && github.event.pull_request.html_url || (github.event_name == 'issue_comment' && github.event.issue.pull_request.html_url) || ''}}  
+      input_pr_title: ${{github.event_name == 'pull_request' && github.event.pull_request.title || (github.event_name == 'issue_comment' && github.event.issue.pull_request.title) || ''}}  
+      repository_name: ${{ github.repository }}  
+      base_ref: ${{github.event_name == 'pull_request' && github.event.pull_request.base.ref || (github.event_name == 'issue_comment' && github.event.issue.pull_request.base.ref) || github.ref}}  
+      trigger_event_type: ${{ github.event_name }}  
+      comment_text: ${{ github.event_name == 'issue_comment' && github.event.comment.body || '' }}

projects/hipother/.github/workflows/validate-pr-description.yml

@@ -0,0 +1,22 @@
+name: Validate PR desription
+
+on:
+  pull_request:
+    types: [opened, edited, synchronize]
+
+jobs:
+  validate-pr-description:
+    runs-on: AMD-ROCm-Internal-dev1
+    steps:
+      - name: Checkout repository
+        uses: actions/checkout@v4
+
+      - name: Set up Python
+        uses: actions/setup-python@v5
+        with:
+          python-version: "3.13"
+
+      - name: Validate PR description
+        env:
+          PR_DESCRIPTION: ${{ github.event.pull_request.body }}
+        run: python .github/scripts/validate_pr_description.py

projects/hipother/CODEOWNERS

@@ -0,0 +1,2 @@
+* @cpaquot_amdeng @gandryey_amdeng @skudchad_amdeng @lmoriche_amdeng
+

projects/hipother/CONTRIBUTING.md

@@ -0,0 +1,133 @@
+# Contributing to hipother #
+
+We welcome contributions to the hipother project. Please follow these details to help ensure your contributions will be successfully accepted.
+If you want to contribute to our documentation, refer to {doc}`Contribute to ROCm docs <rocm:contribute/contributing>`.
+
+## Issue Discussion ##
+
+Please use the [GitHub Issue](https://github.com/ROCm/hipother/issues) tab to notify us of issues.
+
+* Use your best judgement for issue creation. If your issue is already listed, upvote the issue and
+  comment or post to provide additional details, such as how you reproduced this issue.
+* If you're not sure if your issue is the same, err on the side of caution and file your issue.
+  You can add a comment to include the issue number (and link) for the similar issue. If we evaluate
+  your issue as being the same as the existing issue, we'll close the duplicate.
+* If your issue doesn't exist, use the issue template to file a new issue.
+  * When filing an issue, be sure to provide as much information as possible, including script output so
+    we can collect information about your configuration. This helps reduce the time required to
+    reproduce your issue.
+  * Check your issue regularly, as we may require additional information to successfully reproduce the
+    issue.
+* You may also open an issue to ask questions to the maintainers about whether a proposed change
+  meets the acceptance criteria, or to discuss an idea pertaining to the library.
+
+## Acceptance Criteria ##
+
+HIPOTHER is a C++ Runtime API interface with CUDA APIs that allows developers to create portable applications for AMD and NVIDIA GPUs from single source code. Contributors wishing to submit new HIP Features (ie functions, classes, types) should also consider CUDA APIs.
+Differences or limitations of HIP APIs as compared to CUDA APIs should be clearly documented and described.
+Some guidelines are outlined below:
+
+### Add a new HIP API ###
+
+- Add a translation to the hipify-clang tool ; many examples abound.
+    - For stat tracking purposes, place the API into an appropriate stat category ("dev", "mem", "stream", etc).
+- Add a inlined NVIDIA implementation for the function in /hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h in the repository [hipother](https://github.com/ROCm/hipother).
+    - These are typically headers
+- Add an HIP definition and Doxygen comments for the function in /include/hip/hip_runtime_api.h, in the repository [hip](https://github.com/ROCm/hip).
+    - Source implementation typically go in clr/hipamd/src/hip_*.cpp in the reposotory [clr](https://github.com/ROCm/clr). The implementation involves calls to HIP runtime (ie for hipStream_t).
+
+### Run Unit Tests ###
+
+For new features or bug fixes, it's mandatory to run associate [hip-tests](https://github.com/ROCm/hip-tests) on both AMD and NVIDIA platforms.
+Please go to the repo and follow the steps.
+
+For applications and benchmarks outside the hip-tests environment, developments should use a two-step development flow:
+- #1. Compile, link, and install HIP. See {ref}`Building the HIP runtime` notes.
+- #2. Relink the target application to include changes in HIP runtime file.
+
+## Code Structure ##
+
+hipother contains mainly header files with interfaces of different typs of HIP APIs to the corresponding CUDA runtime or driver APIs, for example,
+- `hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h` - contains enumerations and HIP runtime API interfaces corresponding to CUDA enumerations and runtime APIs.
+
+- `hipnv/include/hip/nvidia_detail/nvidia_hiprtc.h` - contains HIP runtime compiler enumerations and APIs correspond to CUDA.
+
+
+## Coding Style ##
+- Code Indentation:
+    - Tabs should be expanded to spaces.
+    - Use 4 spaces indentation.
+- Capitalization and Naming
+    - Prefer camelCase for HIP interfaces and internal symbols.  Note HCC uses _ for separator.
+    - Member variables should begin with a leading "_".  This allows them to be easily distinguished from other variables or functions.
+
+- `{}` placement
+    - namespace should be on same line as `{` and separated by a space.
+    - Single-line if statement should still use `{/}` pair (even though C++ does not require).
+    - For functions, the opening `{` should be placed on a new line.
+    - For if/else blocks, the opening `{` is placed on same line as the if/else. Use a space to separate `{` from if/else. For example,
+```console
+    if (foo) {
+        doFoo()
+    } else {
+        doFooElse();
+    }
+```
+
+
+## Pull Request Guidelines ##
+
+By creating a pull request, you agree to the statements made in the code license section. Your pull request should target the default branch. Our current default branch is the develop branch, which serves as our integration branch.
+
+Follow existing best practice for writing a good Git commit message.
+
+Some tips:
+    http://chris.beams.io/posts/git-commit/
+    https://robots.thoughtbot.com/5-useful-tips-for-a-better-commit-message
+
+In particular :
+   - Use imperative voice, ie "Fix this bug", "Refactor the XYZ routine", "Update the doc".
+     Not : "Fixing the bug", "Fixed the bug", "Bug fix", etc.
+   - Subject should summarize the commit.  Do not end subject with a period.  Use a blank line
+     after the subject.
+
+### Deliverables ###
+
+hipother is part of HIP open source library. Because of this, we include the following license description at the top of every source file.
+If you create new source files in the repository, please include this text in them as well (replacing "xx" with the digits for the current year):
+```
+// Copyright (c) 20xx 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.
+```
+
+### Process ###
+
+After you create a PR, you can take a look at a diff of the changes you made using the PR's "Files" tab.
+
+PRs must pass through the checks and the code review described in the [Acceptance Criteria](#acceptance-criteria) section before they can be merged.
+
+Checks may take some time to complete. You can view their progress in the table near the bottom of the pull request page. You may also be able to use the links in the table
+to view logs associated with a check if it fails.
+
+During code reviews, another developer will take a look through your proposed change. If any modifications are requested (or further discussion about anything is
+needed), they may leave a comment. You can follow up and respond to the comment, and/or create comments of your own if you have questions or ideas.
+When a modification request has been completed, the conversation thread about it will be marked as resolved.
+
+To update the code in your PR (eg. in response to a code review discussion), you can simply push another commit to the branch used in your pull request.

projects/hipother/LICENSE.txt

@@ -0,0 +1,20 @@
+Copyright (c) 2008 - 2024 Advanced Micro Devices, Inc.
+
+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.
+

projects/hipother/README.md

@@ -0,0 +1,22 @@
+## What is this repository for? ###
+
+This repository provides files required to support non-AMD specific back-end implementation for [HIP](https://github.com/ROCm/HIP).
+
+## DISCLAIMER
+
+The information presented in this document is for informational purposes only and may contain technical inaccuracies, omissions, and typographical errors. The information contained herein is subject to change and may be rendered inaccurate for many reasons, including but not limited to product and roadmap changes, component and motherboard versionchanges, new model and/or product releases, product differences between differing manufacturers, software changes, BIOS flashes, firmware upgrades, or the like. Any computer system has risks of security vulnerabilities that cannot be completely prevented or mitigated.AMD assumes no obligation to update or otherwise correct or revise this information. However, AMD reserves the right to revise this information and to make changes from time to time to the content hereof without obligation of AMD to notify any person of such revisions or changes.THIS INFORMATION IS PROVIDED ‘AS IS.” AMD MAKES NO REPRESENTATIONS OR WARRANTIES WITH RESPECT TO THE CONTENTS HEREOF AND ASSUMES NO RESPONSIBILITY FOR ANY INACCURACIES, ERRORS, OR OMISSIONS THAT MAY APPEAR IN THIS INFORMATION. AMD SPECIFICALLY DISCLAIMS ANY IMPLIED WARRANTIES OF NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR ANY PARTICULAR PURPOSE. IN NO EVENT WILL AMD BE LIABLE TO ANY PERSON FOR ANY RELIANCE, DIRECT, INDIRECT, SPECIAL, OR OTHER CONSEQUENTIAL DAMAGES ARISING FROM THE USE OF ANY INFORMATION CONTAINED HEREIN, EVEN IF AMD IS EXPRESSLY ADVISED OF THE POSSIBILITY OF SUCH DAMAGES. AMD, the AMD Arrow logo, and combinations thereof are trademarks of Advanced Micro Devices, Inc. Other product names used in this publication are for identification purposes only and may be trademarks of their respective companies.
+
+©2023 Advanced Micro Devices, Inc. All Rights Reserved.
+
+## Repository branches:
+
+The hipother repository maintains several branches. The branches that are of importance are:
+
+* Develop branch: This is the default branch, on which the new features are still under development and visible. While this maybe of interest to many, it should be noted that this branch and the features under development might not be stable.
+* Release branches. These are branches corresponding to each ROCM release, listed with release tags, such as rocm-6.0, etc.
+
+## Release tagging:
+
+hipother releases are typically naming convention for each ROCM release to help differentiate them.
+
+* rocm x.yy: These are the stable releases based on the ROCM release.

projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_channel_descriptor.h

@@ -0,0 +1,28 @@
+/*
+Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#ifndef HIP_INCLUDE_HIP_NVIDIA_DETAIL_CHANNEL_DESCRIPTOR_H
+#define HIP_INCLUDE_HIP_NVIDIA_DETAIL_CHANNEL_DESCRIPTOR_H
+
+#include "channel_descriptor.h"
+
+#endif

projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_atomics.h

@@ -0,0 +1,67 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#ifndef HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_ATOMICS_H
+#define HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_ATOMICS_H
+
+
+__device__ inline float atomicMax(float* addr, float val) {
+    int ret = __float_as_int(*addr);
+    while (val > __int_as_float(ret)) {
+        int old = ret;
+        if ((ret = atomicCAS((int *)addr, old, __float_as_int(val))) == old)
+            break;
+    }
+    return __int_as_float(ret);
+}
+__device__ inline double atomicMax(double* addr, double val) {
+    unsigned long long ret = __double_as_longlong(*addr);
+    while (val > __longlong_as_double(ret)) {
+        unsigned long long old = ret;
+        if ((ret = atomicCAS((unsigned long long *)addr, old, __double_as_longlong(val))) == old)
+            break;
+    }
+    return __longlong_as_double(ret);
+}
+
+__device__ inline float atomicMin(float* addr, float val) {
+    int ret = __float_as_int(*addr);
+    while (val < __int_as_float(ret)) {
+        int old = ret;
+        if ((ret = atomicCAS((int *)addr, old, __float_as_int(val))) == old)
+            break;
+    }
+    return __int_as_float(ret);
+}
+
+__device__ inline double atomicMin(double* addr, double val) {
+    unsigned long long ret = __double_as_longlong(*addr);
+    while (val < __longlong_as_double(ret)) {
+        unsigned long long old = ret;
+        if ((ret = atomicCAS((unsigned long long *)addr, old, __double_as_longlong(val))) == old)
+            break;
+    }
+    return __longlong_as_double(ret);
+}
+
+
+#endif

projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_bf16.h

@@ -0,0 +1,39 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#ifndef HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_FP16_H
+#define HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_FP16_H
+
+#define HIPRT_ONE_BF16 CUDART_ONE_BF16
+#define HIPRT_ZERO_BF16 CUDART_ZERO_BF16
+#define HIPRT_INF_BF16 CUDART_INF_BF16
+#define HIPRT_MAX_NORMAL_BF16 CUDART_MAX_NORMAL_BF16
+#define HIPRT_MIN_DENORM_BF16 CUDART_MIN_DENORM_BF16
+#define HIPRT_NAN_BF16 CUDART_NAN_BF16
+#define HIPRT_NEG_ZERO_BF16 CUDART_NEG_ZERO_BF16
+
+#include <cuda_bf16.h>
+
+typedef struct __nv_bfloat16 __hip_bfloat16;
+typedef struct __nv_bfloat162 __hip_bfloat162;
+
+#endif // HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_FP16_H

projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_complex.h

@@ -0,0 +1,119 @@
+/*
+Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#ifndef HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_COMPLEX_H
+#define HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_COMPLEX_H
+
+#include "cuComplex.h"
+
+typedef cuFloatComplex hipFloatComplex;
+
+__device__ __host__ static inline float hipCrealf(hipFloatComplex z) { return cuCrealf(z); }
+
+__device__ __host__ static inline float hipCimagf(hipFloatComplex z) { return cuCimagf(z); }
+
+__device__ __host__ static inline hipFloatComplex make_hipFloatComplex(float a, float b) {
+    return make_cuFloatComplex(a, b);
+}
+
+__device__ __host__ static inline hipFloatComplex hipConjf(hipFloatComplex z) { return cuConjf(z); }
+
+__device__ __host__ static inline float hipCsqabsf(hipFloatComplex z) {
+    return cuCabsf(z) * cuCabsf(z);
+}
+
+__device__ __host__ static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q) {
+    return cuCaddf(p, q);
+}
+
+__device__ __host__ static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q) {
+    return cuCsubf(p, q);
+}
+
+__device__ __host__ static inline hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q) {
+    return cuCmulf(p, q);
+}
+
+__device__ __host__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q) {
+    return cuCdivf(p, q);
+}
+
+__device__ __host__ static inline float hipCabsf(hipFloatComplex z) { return cuCabsf(z); }
+
+typedef cuDoubleComplex hipDoubleComplex;
+
+__device__ __host__ static inline double hipCreal(hipDoubleComplex z) { return cuCreal(z); }
+
+__device__ __host__ static inline double hipCimag(hipDoubleComplex z) { return cuCimag(z); }
+
+__device__ __host__ static inline hipDoubleComplex make_hipDoubleComplex(double a, double b) {
+    return make_cuDoubleComplex(a, b);
+}
+
+__device__ __host__ static inline hipDoubleComplex hipConj(hipDoubleComplex z) { return cuConj(z); }
+
+__device__ __host__ static inline double hipCsqabs(hipDoubleComplex z) {
+    return cuCabs(z) * cuCabs(z);
+}
+
+__device__ __host__ static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q) {
+    return cuCadd(p, q);
+}
+
+__device__ __host__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q) {
+    return cuCsub(p, q);
+}
+
+__device__ __host__ static inline hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q) {
+    return cuCmul(p, q);
+}
+
+__device__ __host__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q) {
+    return cuCdiv(p, q);
+}
+
+__device__ __host__ static inline double hipCabs(hipDoubleComplex z) { return cuCabs(z); }
+
+typedef cuFloatComplex hipComplex;
+
+__device__ __host__ static inline hipComplex make_hipComplex(float x, float y) {
+    return make_cuComplex(x, y);
+}
+
+__device__ __host__ static inline hipFloatComplex hipComplexDoubleToFloat(hipDoubleComplex z) {
+    return cuComplexDoubleToFloat(z);
+}
+
+__device__ __host__ static inline hipDoubleComplex hipComplexFloatToDouble(hipFloatComplex z) {
+    return cuComplexFloatToDouble(z);
+}
+
+__device__ __host__ static inline hipComplex hipCfmaf(hipComplex p, hipComplex q, hipComplex r) {
+    return cuCfmaf(p, q, r);
+}
+
+__device__ __host__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, hipDoubleComplex q,
+                                                           hipDoubleComplex r) {
+    return cuCfma(p, q, r);
+}
+
+#endif

projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_cooperative_groups.h

@@ -0,0 +1,12 @@
+#ifndef HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_COOPERATIVE_GROUPS_H
+#define HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_COOPERATIVE_GROUPS_H
+
+// Include CUDA headers
+#include <cuda_runtime.h>
+#include <cooperative_groups.h>
+
+// Include HIP wrapper headers around CUDA
+#include <hip/hip_runtime.h>
+#include <hip/hip_runtime_api.h>
+
+#endif // HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_COOPERATIVE_GROUPS_H

projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_gl_interop.h

@@ -0,0 +1,44 @@
+/*
+Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+#ifndef HIP_INCLUDE_NVIDIA_HIP_GL_INTEROP_H
+#define HIP_INCLUDE_NVIDIA_HIP_GL_INTEROP_H
+
+#include <cuda_gl_interop.h>
+
+typedef enum cudaGLDeviceList hipGLDeviceList;
+#define hipGLDeviceListAll cudaGLDeviceListAll
+#define hipGLDeviceListCurrentFrame  cudaGLDeviceListCurrentFrame
+#define hipGLDeviceListNextFrame  cudaGLDeviceListNextFrame
+
+inline static hipError_t hipGLGetDevices(unsigned int* pHipDeviceCount, int* pHipDevices, unsigned int hipDeviceCount,
+                                         hipGLDeviceList deviceList) {
+  return hipCUDAErrorTohipError(cudaGLGetDevices(pHipDeviceCount, pHipDevices, hipDeviceCount, deviceList));
+}
+
+inline static hipError_t hipGraphicsGLRegisterBuffer(hipGraphicsResource** resource, GLuint buffer, unsigned int flags) {
+  return hipCUDAErrorTohipError(cudaGraphicsGLRegisterBuffer(resource, buffer, flags));
+}
+
+inline static hipError_t hipGraphicsGLRegisterImage(hipGraphicsResource** resource, GLuint image, GLenum target, unsigned int flags) {
+  return hipCUDAErrorTohipError(cudaGraphicsGLRegisterImage(resource, image, target, flags));
+}
+#endif

projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_math_constants.h

@@ -0,0 +1,126 @@
+/*
+Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+#ifndef NVIDIA_HIP_MATH_CONSTANTS_H
+#define NVIDIA_HIP_MATH_CONSTANTS_H
+
+#include <math_constants.h>
+
+// single precision constants
+#define HIP_INF_F            CUDART_INF_F
+#define HIP_NAN_F            CUDART_NAN_F
+#define HIP_MIN_DENORM_F     CUDART_MIN_DENORM_F
+#define HIP_MAX_NORMAL_F     CUDART_MAX_NORMAL_F
+#define HIP_NEG_ZERO_F       CUDART_NEG_ZERO_F
+#define HIP_ZERO_F           CUDART_ZERO_F
+#define HIP_ONE_F            CUDART_ONE_F
+#define HIP_SQRT_HALF_F      CUDART_SQRT_HALF_F
+#define HIP_SQRT_HALF_HI_F   CUDART_SQRT_HALF_HI_F
+#define HIP_SQRT_HALF_LO_F   CUDART_SQRT_HALF_LO_F
+#define HIP_SQRT_TWO_F       CUDART_SQRT_TWO_F
+#define HIP_THIRD_F          CUDART_THIRD_F
+#define HIP_PIO4_F           CUDART_PIO4_F
+#define HIP_PIO2_F           CUDART_PIO2_F
+#define HIP_3PIO4_F          CUDART_3PIO4_F
+#define HIP_2_OVER_PI_F      CUDART_2_OVER_PI_F
+#define HIP_SQRT_2_OVER_PI_F CUDART_SQRT_2_OVER_PI_F
+#define HIP_PI_F             CUDART_PI_F
+#define HIP_L2E_F            CUDART_L2E_F
+#define HIP_L2T_F            CUDART_L2T_F
+#define HIP_LG2_F            CUDART_LG2_F
+#define HIP_LGE_F            CUDART_LGE_F
+#define HIP_LN2_F            CUDART_LN2_F
+#define HIP_LNT_F            CUDART_LNT_F
+#define HIP_LNPI_F           CUDART_LNPI_F
+#define HIP_TWO_TO_M126_F    CUDART_TWO_TO_M126_F
+#define HIP_TWO_TO_126_F     CUDART_TWO_TO_126_F
+#define HIP_NORM_HUGE_F      CUDART_NORM_HUGE_F
+#define HIP_TWO_TO_23_F      CUDART_TWO_TO_23_F
+#define HIP_TWO_TO_24_F      CUDART_TWO_TO_24_F
+#define HIP_TWO_TO_31_F      CUDART_TWO_TO_31_F
+#define HIP_TWO_TO_32_F      CUDART_TWO_TO_32_F
+#define HIP_REMQUO_BITS_F    CUDART_REMQUO_BITS_F
+#define HIP_REMQUO_MASK_F    CUDART_REMQUO_MASK_F
+#define HIP_TRIG_PLOSS_F     CUDART_TRIG_PLOSS_F
+
+// double precision constants
+#define HIP_INF              CUDART_INF
+#define HIP_NAN              CUDART_NAN
+#define HIP_NEG_ZERO         CUDART_NEG_ZERO
+#define HIP_MIN_DENORM       CUDART_MIN_DENORM
+#define HIP_ZERO             CUDART_ZERO
+#define HIP_ONE              CUDART_ONE
+#define HIP_SQRT_TWO         CUDART_SQRT_TWO
+#define HIP_SQRT_HALF        CUDART_SQRT_HALF
+#define HIP_SQRT_HALF_HI     CUDART_SQRT_HALF_HI
+#define HIP_SQRT_HALF_LO     CUDART_SQRT_HALF_LO
+#define HIP_THIRD            CUDART_THIRD
+#define HIP_TWOTHIRD         CUDART_TWOTHIRD
+#define HIP_PIO4             CUDART_PIO4
+#define HIP_PIO4_HI          CUDART_PIO4_HI
+#define HIP_PIO4_LO          CUDART_PIO4_LO
+#define HIP_PIO2             CUDART_PIO2
+#define HIP_PIO2_HI          CUDART_PIO2_HI
+#define HIP_PIO2_LO          CUDART_PIO2_LO
+#define HIP_3PIO4            CUDART_3PIO4
+#define HIP_2_OVER_PI        CUDART_2_OVER_PI
+#define HIP_PI               CUDART_PI
+#define HIP_PI_HI            CUDART_PI_HI
+#define HIP_PI_LO            CUDART_PI_LO
+#define HIP_SQRT_2PI         CUDART_SQRT_2PI
+#define HIP_SQRT_2PI_HI      CUDART_SQRT_2PI_HI
+#define HIP_SQRT_2PI_LO      CUDART_SQRT_2PI_LO
+#define HIP_SQRT_PIO2        CUDART_SQRT_PIO2
+#define HIP_SQRT_PIO2_HI     CUDART_SQRT_PIO2_HI
+#define HIP_SQRT_PIO2_LO     CUDART_SQRT_PIO2_LO
+#define HIP_SQRT_2OPI        CUDART_SQRT_2OPI
+#define HIP_L2E              CUDART_L2E
+#define HIP_L2E_HI           CUDART_L2E_HI
+#define HIP_L2E_LO           CUDART_L2E_LO
+#define HIP_L2T              CUDART_L2T
+#define HIP_LG2              CUDART_LG2
+#define HIP_LG2_HI           CUDART_LG2_HI
+#define HIP_LG2_LO           CUDART_LG2_LO
+#define HIP_LGE              CUDART_LGE
+#define HIP_LGE_HI           CUDART_LGE_HI
+#define HIP_LGE_LO           CUDART_LGE_LO
+#define HIP_LN2              CUDART_LN2
+#define HIP_LN2_HI           CUDART_LN2_HI
+#define HIP_LN2_LO           CUDART_LN2_LO
+#define HIP_LNT              CUDART_LNT
+#define HIP_LNT_HI           CUDART_LNT_HI
+#define HIP_LNT_LO           CUDART_LNT_LO
+#define HIP_LNPI             CUDART_LNPI
+#define HIP_LN2_X_1024       CUDART_LN2_X_1024
+#define HIP_LN2_X_1025       CUDART_LN2_X_1025
+#define HIP_LN2_X_1075       CUDART_LN2_X_1075
+#define HIP_LG2_X_1024       CUDART_LG2_X_1024
+#define HIP_LG2_X_1075       CUDART_LG2_X_1075
+#define HIP_TWO_TO_23        CUDART_TWO_TO_23
+#define HIP_TWO_TO_52        CUDART_TWO_TO_52
+#define HIP_TWO_TO_53        CUDART_TWO_TO_53
+#define HIP_TWO_TO_54        CUDART_TWO_TO_54
+#define HIP_TWO_TO_M54       CUDART_TWO_TO_M54
+#define HIP_TWO_TO_M1022     CUDART_TWO_TO_M1022
+#define HIP_TRIG_PLOSS       CUDART_TRIG_PLOSS
+#define HIP_DBL2INT_CVT      CUDART_DBL2INT_CVT
+
+#endif

projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime.h

@@ -0,0 +1,124 @@
+/*
+Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#ifndef HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_RUNTIME_H
+#define HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_RUNTIME_H
+
+#include <cuda_runtime.h>
+
+#include <hip/hip_runtime_api.h>
+
+#define HIP_KERNEL_NAME(...) __VA_ARGS__
+
+typedef int hipLaunchParm;
+
+#define hipLaunchKernelGGLInternal(kernelName, numBlocks, numThreads, memPerBlock, streamId, ...)  \
+  do {                                                                                             \
+    kernelName<<<numBlocks, numThreads, memPerBlock, streamId>>>(__VA_ARGS__);                     \
+  } while (0)
+
+#define hipLaunchKernelGGL(kernelName, ...) hipLaunchKernelGGLInternal((kernelName), __VA_ARGS__)
+
+#define hipReadModeElementType cudaReadModeElementType
+
+#ifdef __CUDA_ARCH__
+
+
+// 32-bit Atomics:
+#define __HIP_ARCH_HAS_GLOBAL_INT32_ATOMICS__ (__CUDA_ARCH__ >= 110)
+#define __HIP_ARCH_HAS_GLOBAL_FLOAT_ATOMIC_EXCH__ (__CUDA_ARCH__ >= 110)
+#define __HIP_ARCH_HAS_SHARED_INT32_ATOMICS__ (__CUDA_ARCH__ >= 120)
+#define __HIP_ARCH_HAS_SHARED_FLOAT_ATOMIC_EXCH__ (__CUDA_ARCH__ >= 120)
+#define __HIP_ARCH_HAS_FLOAT_ATOMIC_ADD__ (__CUDA_ARCH__ >= 200)
+
+// 64-bit Atomics:
+#define __HIP_ARCH_HAS_GLOBAL_INT64_ATOMICS__ (__CUDA_ARCH__ >= 200)
+#define __HIP_ARCH_HAS_SHARED_INT64_ATOMICS__ (__CUDA_ARCH__ >= 120)
+
+// Doubles
+#define __HIP_ARCH_HAS_DOUBLES__ (__CUDA_ARCH__ >= 120)
+
+// warp cross-lane operations:
+#define __HIP_ARCH_HAS_WARP_VOTE__ (__CUDA_ARCH__ >= 120)
+#define __HIP_ARCH_HAS_WARP_BALLOT__ (__CUDA_ARCH__ >= 200)
+#define __HIP_ARCH_HAS_WARP_SHUFFLE__ (__CUDA_ARCH__ >= 300)
+#define __HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ (__CUDA_ARCH__ >= 350)
+
+// sync
+#define __HIP_ARCH_HAS_THREAD_FENCE_SYSTEM__ (__CUDA_ARCH__ >= 200)
+#define __HIP_ARCH_HAS_SYNC_THREAD_EXT__ (__CUDA_ARCH__ >= 200)
+
+// misc
+#define __HIP_ARCH_HAS_SURFACE_FUNCS__ (__CUDA_ARCH__ >= 200)
+#define __HIP_ARCH_HAS_3DGRID__ (__CUDA_ARCH__ >= 200)
+#define __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ (__CUDA_ARCH__ >= 350)
+
+#endif
+
+#ifdef __CUDACC__
+
+#include "nvidia_hip_atomics.h"
+#include "nvidia_hip_unsafe_atomics.h"
+
+#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 HIP_SYMBOL(X) &X
+
+/**
+ * Map HIP_DYNAMIC_SHARED to "extern __shared__" for compatibility with old HIP applications
+ * To be removed in a future release.
+ */
+#define HIP_DYNAMIC_SHARED(type, var) extern __shared__ type var[];
+#define HIP_DYNAMIC_SHARED_ATTRIBUTE
+
+#ifdef __HIP_DEVICE_COMPILE__
+#define abort_()                                                                                   \
+  { asm("trap;"); }
+#undef assert
+#define assert(COND)                                                                               \
+  {                                                                                                \
+    if (!COND) {                                                                                   \
+      abort_();                                                                                    \
+    }                                                                                              \
+  }
+#endif
+
+#define __clock() clock()
+#define __clock64() clock64()
+
+#endif
+
+#endif

projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_texture_types.h

@@ -0,0 +1,6 @@
+#ifndef HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_TEXTURE_TYPES_H
+#define HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_TEXTURE_TYPES_H
+
+#include <texture_types.h>
+
+#endif

projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_unsafe_atomics.h

@@ -0,0 +1,100 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#ifndef HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_UNSAFE_ATOMICS_H
+#define HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_UNSAFE_ATOMICS_H
+
+__device__ inline float unsafeAtomicAdd(float* addr, float value) {
+    return atomicAdd(addr, value);
+}
+
+__device__ inline double unsafeAtomicAdd(double* addr, double value) {
+#if __CUDA_ARCH__ < 600
+    unsigned long long *addr_cast = (unsigned long long*)addr;
+    unsigned long long old_val = *addr_cast;
+    unsigned long long expected;
+    do {
+        expected = old_val;
+        old_val = atomicCAS(addr_cast, expected,
+                            __double_as_longlong(value +
+                                                 __longlong_as_double(expected)));
+    } while (__double_as_longlong(expected) != __double_as_longlong(old_val));
+    return old_val;
+#else
+    return atomicAdd(addr, value);
+#endif
+}
+
+__device__ inline float unsafeAtomicMax(float* addr, float value) {
+    return atomicMax(addr, value);
+}
+
+__device__ inline double unsafeAtomicMax(double* addr, double val) {
+    return atomicMax(addr, val);
+}
+
+__device__ inline float unsafeAtomicMin(float* addr, float value) {
+    return atomicMin(addr, value);
+}
+
+__device__ inline double unsafeAtomicMin(double* addr, double val) {
+    return atomicMin(addr, val);
+}
+
+__device__ inline float safeAtomicAdd(float* addr, float value) {
+    return atomicAdd(addr, value);
+}
+
+__device__ inline double safeAtomicAdd(double* addr, double value) {
+#if __CUDA_ARCH__ < 600
+    unsigned long long *addr_cast = (unsigned long long*)addr;
+    unsigned long long old_val = *addr_cast;
+    unsigned long long expected;
+    do {
+        expected = old_val;
+        old_val = atomicCAS(addr_cast, expected,
+                            __double_as_longlong(value +
+                                                 __longlong_as_double(expected)));
+    } while (__double_as_longlong(expected) != __double_as_longlong(old_val));
+    return old_val;
+#else
+    return atomicAdd(addr, value);
+#endif
+}
+
+__device__ inline float safeAtomicMax(float* addr, float value) {
+    return atomicMax(addr, value);
+}
+
+__device__ inline double safeAtomicMax(double* addr, double val) {
+    return atomicMax(addr, val);
+}
+
+__device__ inline float safeAtomicMin(float* addr, float value) {
+    return atomicMin(addr, value);
+}
+
+__device__ inline double safeAtomicMin(double* addr, double val) {
+    return atomicMin(addr, val);
+}
+
+#endif

projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hiprtc.h

@@ -0,0 +1,172 @@
+/*
+Copyright (c) 2021 - 2022 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+#ifndef HIPRTC_H
+#define HIPRTC_H
+
+#include <cuda.h>
+#include <nvrtc.h>
+
+#ifdef __cplusplus
+extern "C" {
+#endif /* __cplusplus */
+
+#include <stdlib.h>
+
+#if !defined(_WIN32)
+#pragma GCC visibility push(default)
+#endif
+
+typedef enum hiprtcResult {
+  HIPRTC_SUCCESS = 0,
+  HIPRTC_ERROR_OUT_OF_MEMORY = 1,
+  HIPRTC_ERROR_PROGRAM_CREATION_FAILURE = 2,
+  HIPRTC_ERROR_INVALID_INPUT = 3,
+  HIPRTC_ERROR_INVALID_PROGRAM = 4,
+  HIPRTC_ERROR_INVALID_OPTION = 5,
+  HIPRTC_ERROR_COMPILATION = 6,
+  HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE = 7,
+  HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = 8,
+  HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = 9,
+  HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID = 10,
+  HIPRTC_ERROR_INTERNAL_ERROR = 11
+} hiprtcResult;
+
+inline static nvrtcResult hiprtcResultTonvrtcResult(hiprtcResult result) {
+  switch (result) {
+    case HIPRTC_SUCCESS:
+      return NVRTC_SUCCESS;
+    case HIPRTC_ERROR_OUT_OF_MEMORY:
+      return NVRTC_ERROR_OUT_OF_MEMORY;
+    case HIPRTC_ERROR_PROGRAM_CREATION_FAILURE:
+      return NVRTC_ERROR_PROGRAM_CREATION_FAILURE;
+    case HIPRTC_ERROR_INVALID_INPUT:
+      return NVRTC_ERROR_INVALID_INPUT;
+    case HIPRTC_ERROR_INVALID_PROGRAM:
+      return NVRTC_ERROR_INVALID_PROGRAM;
+    case HIPRTC_ERROR_INVALID_OPTION:
+      return NVRTC_ERROR_INVALID_OPTION;
+    case HIPRTC_ERROR_COMPILATION:
+      return NVRTC_ERROR_COMPILATION;
+    case HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE:
+      return NVRTC_ERROR_BUILTIN_OPERATION_FAILURE;
+    case HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION:
+      return NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION;
+    case HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION:
+      return NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION;
+    case HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID:
+      return NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID;
+    case HIPRTC_ERROR_INTERNAL_ERROR:
+      return NVRTC_ERROR_INTERNAL_ERROR;
+    default:
+      return NVRTC_ERROR_INTERNAL_ERROR;
+  }
+}
+
+inline static hiprtcResult nvrtcResultTohiprtcResult(nvrtcResult result) {
+  switch (result) {
+    case NVRTC_SUCCESS:
+      return HIPRTC_SUCCESS;
+    case NVRTC_ERROR_OUT_OF_MEMORY:
+      return HIPRTC_ERROR_OUT_OF_MEMORY;
+    case NVRTC_ERROR_PROGRAM_CREATION_FAILURE:
+      return HIPRTC_ERROR_PROGRAM_CREATION_FAILURE;
+    case NVRTC_ERROR_INVALID_INPUT:
+      return HIPRTC_ERROR_INVALID_INPUT;
+    case NVRTC_ERROR_INVALID_PROGRAM:
+      return HIPRTC_ERROR_INVALID_PROGRAM;
+    case NVRTC_ERROR_INVALID_OPTION:
+      return HIPRTC_ERROR_INVALID_OPTION;
+    case NVRTC_ERROR_COMPILATION:
+      return HIPRTC_ERROR_COMPILATION;
+    case NVRTC_ERROR_BUILTIN_OPERATION_FAILURE:
+      return HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE;
+    case NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION:
+      return HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION;
+    case NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION:
+      return HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION;
+    case NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID:
+      return HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID;
+    case NVRTC_ERROR_INTERNAL_ERROR:
+      return HIPRTC_ERROR_INTERNAL_ERROR;
+    default:
+      return HIPRTC_ERROR_INTERNAL_ERROR;
+  }
+}
+
+inline static const char* hiprtcGetErrorString(hiprtcResult result) {
+  return nvrtcGetErrorString(hiprtcResultTonvrtcResult(result));
+}
+
+inline static hiprtcResult hiprtcVersion(int* major, int* minor) {
+  return nvrtcResultTohiprtcResult(nvrtcVersion(major, minor));
+}
+
+typedef nvrtcProgram hiprtcProgram;
+
+inline static hiprtcResult hiprtcAddNameExpression(hiprtcProgram prog, const char* name_expression) {
+  return nvrtcResultTohiprtcResult(nvrtcAddNameExpression(prog, name_expression));
+}
+
+inline static hiprtcResult hiprtcCompileProgram(hiprtcProgram prog, int numOptions, const char** options) {
+  return nvrtcResultTohiprtcResult(nvrtcCompileProgram(prog, numOptions, options));
+}
+
+inline static hiprtcResult hiprtcCreateProgram(hiprtcProgram* prog, const char* src, const char* name,
+                                 int numHeaders, const char** headers, const char** includeNames) {
+  return nvrtcResultTohiprtcResult(
+      nvrtcCreateProgram(prog, src, name, numHeaders, headers, includeNames));
+}
+
+inline static hiprtcResult hiprtcDestroyProgram(hiprtcProgram* prog) {
+  return nvrtcResultTohiprtcResult(nvrtcDestroyProgram(prog));
+}
+
+inline static hiprtcResult hiprtcGetLoweredName(hiprtcProgram prog, const char* name_expression,
+                                  const char** lowered_name) {
+  return nvrtcResultTohiprtcResult(nvrtcGetLoweredName(prog, name_expression, lowered_name));
+}
+
+inline static hiprtcResult hiprtcGetProgramLog(hiprtcProgram prog, char* log) {
+  return nvrtcResultTohiprtcResult(nvrtcGetProgramLog(prog, log));
+}
+
+inline static hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram prog, size_t* logSizeRet) {
+  return nvrtcResultTohiprtcResult(nvrtcGetProgramLogSize(prog, logSizeRet));
+}
+
+inline static hiprtcResult hiprtcGetCode(hiprtcProgram prog, char* code) {
+  return nvrtcResultTohiprtcResult(nvrtcGetPTX(prog, code));
+}
+
+inline static hiprtcResult hiprtcGetCodeSize(hiprtcProgram prog, size_t* codeSizeRet) {
+  return nvrtcResultTohiprtcResult(nvrtcGetPTXSize(prog, codeSizeRet));
+}
+
+#if !defined(_WIN32)
+#pragma GCC visibility pop
+#endif
+
+#ifdef __cplusplus
+}
+#endif /* __cplusplus */
+
+#endif  // HIPRTC_H