Develop Stream 2024-05-06 implementation (#115)

Develop Stream 2024-05-06 implementation * Add copyright check script based on hipCUB * Implemented minimal CMake CUDA example * Add a default case for externally controlled enumerations * ci: use build instead rocm-build and nvcc-build tags * Add consistent error code for parsing failures * ci: manually set HIP_COMPILER, HIP_PLATFORM and HIP_RUNTIME * Add git-clang-format check when installing pre-commit hook * Fix markdown linting --------- Co-authored-by: Gergely Meszaros <[email protected]> Co-authored-by: Lőrinc Serfőző <[email protected]> Co-authored-by: Jaap <[email protected]> Co-authored-by: Nara Prasetya <[email protected]>
ROCm · Jun 3, 2024 · fd90d0e · fd90d0e
1 parent c707559
commit fd90d0e
Show file tree

Hide file tree

Showing 31 changed files with 487 additions and 77 deletions.
diff --git a/.githooks/install b/.githooks/install
@@ -0,0 +1,14 @@
+#!/bin/sh
+
+cd "$(git rev-parse --git-dir)"
+cd hooks
+
+echo "Installing hooks..."
+# Install pre-commit hook if dependencies are satisfied
+if ! [ -x "$(command -v git-clang-format)" ]; then
+  echo 'Error: pre-commit hook depends on git-clang-format, but is not installed.' >&2
+  exit 1
+else
+    ln -s ../../.githooks/pre-commit pre-commit
+fi
+echo "Done!"
diff --git a/.githooks/pre-commit b/.githooks/pre-commit
@@ -0,0 +1,31 @@
+#!/bin/sh
+
+# Redirect output to stderr.
+exec 1>&2
+
+check_failed=false
+
+# Do the code format check
+if ! "$(git rev-parse --show-toplevel)/Scripts/CodeFormat/check_format.sh" HEAD --cached 1>&2; then
+    printf "\n\033[31mFailed\033[0m: code format check.\n"
+    check_failed=true
+fi
+
+# Do the copyright check
+# update & apply copyright when hook config is set, otherwise just verify
+opts="-qc"
+if [ "$(git config --get --type bool --default false hooks.updateCopyright)" = "true" ]; then
+    opts="-qca"
+fi
+
+if ! "$(git rev-parse --show-toplevel)/Scripts/CopyrightDate/check_copyright.sh" "$opts" 1>&2; then
+    printf "\n\033[31mFailed\033[0m: copyright date check.\n"
+    check_failed=true
+fi
+
+if $check_failed; then
+    printf "
+Pre-commit check failed, please fix the reported errors.
+Note: Use '\033[33mgit commit --no-verify\033[0m' to bypass checks.\n"
+    exit 1
+fi
diff --git a/.gitlab-ci.yml b/.gitlab-ci.yml
@@ -48,7 +48,7 @@ clang-format:
   stage: lint
   needs: []
   tags:
-    - rocm-build
+    - build
   variables:
     CLANG_FORMAT: "/opt/rocm/llvm/bin/clang-format"
     GIT_CLANG_FORMAT: "/opt/rocm/llvm/bin/git-clang-format"
@@ -59,6 +59,19 @@ clang-format:
     - git config --global --add safe.directory $CI_PROJECT_DIR
     - Scripts/CodeFormat/check_format.sh $CI_MERGE_REQUEST_DIFF_BASE_SHA --binary "$CLANG_FORMAT"
 
+copyright-date:
+  image: $DOCKER_TAG_PREFIX:rocm-ubuntu
+  stage: lint
+  needs: []
+  tags:
+    - build
+  rules:
+    - if: '$CI_PIPELINE_SOURCE == "merge_request_event"'
+  script:
+    - cd $CI_PROJECT_DIR
+    - git config --global --add safe.directory $CI_PROJECT_DIR
+    - Scripts/CopyrightDate/check_copyright.sh -v -d $CI_MERGE_REQUEST_DIFF_BASE_SHA
+
 .build:dockerfiles:
   timeout: 60m
   image:
@@ -67,7 +80,7 @@ clang-format:
   stage: build
   needs: []
   tags:
-    - rocm-build
+    - build
   script:
     - mkdir -p /kaniko/.docker
     - echo "${DOCKER_AUTH_CONFIG}" > /kaniko/.docker/config.json
@@ -108,7 +121,7 @@ build:make-rocm:
   extends:
     - .rules:build
   tags:
-    - rocm-build
+    - build
   needs: []
   script:
     - cd $CI_PROJECT_DIR && make CXXFLAGS="$HIP_FLAGS" -j $(nproc)
@@ -119,7 +132,7 @@ build:make-cuda:
   extends:
     - .rules:build
   tags:
-    - nvcc-build
+    - build
   needs: []
   script:
     - cd $CI_PROJECT_DIR && make CXXFLAGS="$CUDA_FLAGS" GPU_RUNTIME=CUDA -j $(nproc)
@@ -143,7 +156,7 @@ build:cmake-rocm:
     - .build:cmake
     - .gpus:rocm-gpus
   tags:
-    - rocm-build
+    - build
   script:
     - cmake
       -S $CI_PROJECT_DIR
@@ -166,7 +179,7 @@ build:cmake-cuda:
  extends:
    - .build:cmake
  tags:
-   - nvcc-build
+   - build
  script:
    - cmake
      -S $CI_PROJECT_DIR
@@ -209,6 +222,10 @@ test:cuda:
   extends:
     - .test
     - .gpus:nvcc
+  before_script:
+    - export HIP_COMPILER=nvcc
+    - export HIP_PLATFORM=nvidia
+    - export HIP_RUNTIME=cuda
   needs:
     - build:cmake-cuda
 

diff --git a/Applications/bitonic_sort/main.hip b/Applications/bitonic_sort/main.hip
@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2023-2024 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
@@ -148,7 +148,7 @@ int main(int argc, char* argv[])
     {
         std::cout << "The ordering must be 'dec' or 'inc', the default ordering is 'inc'."
                   << std::endl;
-        return 0;
+        return error_exit_code;
     }
     const bool sort_increasing = (sort.compare("inc") == 0);
 

diff --git a/Applications/monte_carlo_pi/main.hip b/Applications/monte_carlo_pi/main.hip
@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2023-2024 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
@@ -126,7 +126,7 @@ int main(int argc, char* argv[])
     if(sample_count <= 0)
     {
         std::cerr << "Sample count should be greater than 0." << std::endl;
-        return 0;
+        return error_exit_code;
     }
 
     // The samples have two dimensions, so two random numbers are required per sample.

diff --git a/Applications/prefix_sum/main.hip b/Applications/prefix_sum/main.hip
@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2023-2024 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
@@ -182,7 +182,7 @@ int main(int argc, char* argv[])
     if(size <= 0)
     {
         std::cout << "Size must be at least 1." << std::endl;
-        exit(0);
+        return error_exit_code;
     }
 
     // 2. Generate input vector.

diff --git a/Common/hipsolver_utils.hpp b/Common/hipsolver_utils.hpp
@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2023-2024 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
@@ -52,10 +52,10 @@ inline const char* hipsolverStatusToString(hipsolverStatus_t status)
 #if (hipsolverVersionMajor >= 2 && hipsolverVersionMinor >= 1)
         case HIPSOLVER_STATUS_MATRIX_TYPE_NOT_SUPPORTED : return "HIPSOLVER_STATUS_MATRIX_TYPE_NOT_SUPPORTED";
 #endif
+        // We do use default because we are not in control of these enumeration values.
+        // Ideally this function is something hipsolver would provide
+        default: return "<unknown hipsolverStatus_t value>";
     }
-    // We don't use default so that the compiler warns if any valid enums are missing from the
-    // switch. If the value is not a valid hipsolverStatus_t, we return the following.
-    return "<undefined hipsolverStatus_t value>";
 }
 
 /// \brief Checks if the provided status code is \p HIPSOLVER_STATUS_SUCCESS and if not,

diff --git a/Common/rocsparse_utils.hpp b/Common/rocsparse_utils.hpp
@@ -53,10 +53,10 @@ inline const char* rocsparse_status_to_string(rocsparse_status status)
 #endif
         case rocsparse_status_requires_sorted_storage:
             return "rocsparse_status_requires_sorted_storage";
+        // We do use default because we are not in control of these enumeration values.
+        // Ideally this function is something rocsparse would provide
+        default: return "<unknown rocsparse_status value>";
     }
-    // We don't use default so that the compiler warns if any valid enums are missing from the
-    // switch. If the value is not a valid rocsparse_status, we return the following.
-    return "<undefined rocsparse_status value>";
 }
 
 /// \brief Checks if the provided status code is \p rocsparse_status_success and if not,

diff --git a/HIP-Basic/CMakeLists.txt b/HIP-Basic/CMakeLists.txt
@@ -95,6 +95,10 @@ if(NOT WIN32)
         message("Perl not found, not building hipify example")
     endif()
 endif()
+if("${GPU_RUNTIME}" STREQUAL "CUDA")
+    add_subdirectory(hello_world_cuda)
+endif()
+
 add_subdirectory(inline_assembly)
 add_subdirectory(matrix_multiplication)
 add_subdirectory(moving_average)

diff --git a/HIP-Basic/hello_world_cuda/CMakeLists.txt b/HIP-Basic/hello_world_cuda/CMakeLists.txt
@@ -0,0 +1,51 @@
+# MIT License
+#
+# Copyright (c) 2024 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.
+
+set(example_name hip_hello_world_cuda)
+
+cmake_minimum_required(VERSION 3.21 FATAL_ERROR)
+
+# CMake's HIP language mode does not yet support compiling with CUDA, thereby we must
+# resort to the CUDA language mode.
+project(${example_name} LANGUAGES CUDA)
+
+if(WIN32)
+    set(ROCM_ROOT "$ENV{HIP_PATH}" CACHE PATH "Root directory of the ROCm installation")
+else()
+    set(ROCM_ROOT "/opt/rocm" CACHE PATH "Root directory of the ROCm installation")
+endif()
+
+add_executable(${example_name} main.hip)
+
+# Make example runnable using ctest
+add_test(${example_name} ${example_name})
+
+# Make the HIP runtime headers accessible
+target_include_directories(${example_name} PRIVATE
+    "${ROCM_ROOT}/include"
+    "${CMAKE_CURRENT_SOURCE_DIR}/../../Common")
+
+# Set up the compilation language for the source file.
+# Usually this can be deduced from the file extension, but not in the case of .hip.
+set_source_files_properties(main.hip PROPERTIES LANGUAGE CUDA)
+
+install(TARGETS ${example_name})
diff --git a/HIP-Basic/hello_world_cuda/README.md b/HIP-Basic/hello_world_cuda/README.md
@@ -0,0 +1,31 @@
+# HIP-Basic Hello World on the CUDA platform Example
+
+## Description
+
+This example showcases a simple HIP program that is compiled on the CUDA platform using CMake.
+
+### Application flow
+
+1. A kernel is launched: the function `hello_world_kernel` is executed on the device. The kernel is executed on a single thread, and prints "Hello, World!" to the console.
+2. A launch error check is performed using `hipGetLastError`.
+3. _Synchronization_ is performed: the host program execution halts until the kernel on the device has finished executing.
+
+## Key APIs and Concepts
+
+- For introduction to the programming concepts in this example, refer to the general [hello world example](../hello_world/).
+- This example showcases setting up a HIP program to be compiled to the CUDA platform using CMake.
+  - Since CMake (as of version 3.21) does not support compiling to CUDA in HIP language mode, CUDA language mode has to be used. Thereby the project language is specified as `CUDA`.
+  - Additionally, we must "teach" CMake to compile the source file `main.hip` in CUDA language mode, because it cannot guess that from the file extension. This is done by `set_source_files_properties(main.hip PROPERTIES LANGUAGE CUDA)`.
+  - The HIP "runtime" on the CUDA platform is header only. Thereby there is no need to link to a library, but the HIP include directory have to be added to the search paths. This is performed by `target_include_directories(${example_name} PRIVATE "${ROCM_ROOT}/include"`.
+
+## Demonstrated API Calls
+
+### HIP Runtime
+
+- `hipGetLastError`
+- `hipDeviceSynchronize`
+- `__global__`
+
+## Supported Platforms
+
+This example is only supported on the CUDA platform.
diff --git a/HIP-Basic/hello_world_cuda/main.hip b/HIP-Basic/hello_world_cuda/main.hip
@@ -0,0 +1,39 @@
+// MIT License
+//
+// Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to deal
+// in the Software without restriction, including without limitation the rights
+// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the Software is
+// furnished to do so, subject to the following conditions:
+//
+// The above copyright notice and this permission notice shall be included in all
+// copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+// SOFTWARE.
+
+#include <hip/hip_runtime.h>
+
+#include "example_utils.hpp"
+
+__global__ void hello_world_kernel()
+{
+    printf("Hello, World!\n");
+}
+
+int main()
+{
+    static constexpr unsigned int grid_size  = 1;
+    static constexpr unsigned int block_size = 1;
+    hello_world_kernel<<<grid_size, block_size>>>();
+    HIP_CHECK(hipGetLastError());
+    HIP_CHECK(hipDeviceSynchronize());
+}
diff --git a/Libraries/hipBLAS/gemm_strided_batched/main.hip b/Libraries/hipBLAS/gemm_strided_batched/main.hip
@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2022-2023 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2022-204 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
@@ -59,25 +59,25 @@ int main(const int argc, const char** argv)
     if(m <= 0)
     {
         std::cout << "Value of 'm' should be greater than 0" << std::endl;
-        return 0;
+        return error_exit_code;
     }
 
     if(n <= 0)
     {
         std::cout << "Value of 'n' should be greater than 0" << std::endl;
-        return 0;
+        return error_exit_code;
     }
 
     if(k <= 0)
     {
         std::cout << "Value of 'k' should be greater than 0" << std::endl;
-        return 0;
+        return error_exit_code;
     }
 
     if(batch_count <= 0)
     {
         std::cout << "Value of 'c' should be greater than 0" << std::endl;
-        return 0;
+        return error_exit_code;
     }
 
     // Set scalar values used for multiplication.