Add github workflow to compile for HIP and fix some HIP related issues

Author: stijnh

.github/workflows/cmake-action.yml renamed to .github/workflows/cmake-run-cuda.yml

@@ -0,0 +1,48 @@
+name: CMake
+
+on:
+  workflow_call:
+    inputs:
+      rocm-version:
+        required: true
+        type: string
+      gpu-arch:
+        required: false
+        type: string
+        default: "gfx942"
+
+env:
+  # Customize the CMake build type here (Release, Debug, RelWithDebInfo, etc.)
+  BUILD_TYPE: Debug
+
+jobs:
+  build:
+    # The CMake configure and build commands are platform agnostic and should work equally well on Windows or Mac.
+    # You can convert this to a matrix build if you need cross-platform coverage.
+    # See: https://docs.github.com/en/free-pro-team@latest/actions/learn-github-actions/managing-complex-workflows#using-a-build-matrix
+    runs-on: ubuntu-latest
+
+    steps:
+    - uses: loostrum/[email protected]
+      with:
+        version:  ${{ inputs.rocm-version }}
+
+    - uses: actions/checkout@v3
+      with:
+        submodules: 'true'
+
+    - name: Configure CMake
+      # Configure CMake in a 'build' subdirectory. `CMAKE_BUILD_TYPE` is only required if you are using a single-configuration generator such as make.
+      # See https://cmake.org/cmake/help/latest/variable/CMAKE_BUILD_TYPE.html?highlight=cmake_build_type
+      run: cmake -B ${{github.workspace}}/build -DCMAKE_BUILD_TYPE=${{env.BUILD_TYPE}} -DKERNEL_FLOAT_BUILD_TEST=1 -DKERNEL_FLOAT_BUILD_EXAMPLE=1 -DKERNEL_FLOAT_LANGUAGE=HIP -DCMAKE_HIP_ARCHITECTURES=${{ inputs.gpu-arch }}
+
+    - name: Build
+      # Build your program with the given configuration
+      run: cmake --build ${{github.workspace}}/build --config ${{env.BUILD_TYPE}}
+
+    - name: Test
+      working-directory: ${{github.workspace}}/build
+      # Execute tests defined by the CMake configuration.
+      # See https://cmake.org/cmake/help/latest/manual/ctest.1.html for more detail
+      run: ./tests/kernel_float_tests --durations=yes --success --verbosity=high ~[GPU]
+

.github/workflows/cmake-run-hip.yml

@@ -6,23 +6,51 @@ on:
     branches: [ "main" ]
 
 env:
-  # Customize the CMake build type here (Release, Debug, RelWithDebInfo, etc.)
   BUILD_TYPE: Debug
 
 jobs:
   build-cuda:
-    uses: ./.github/workflows/cmake-action.yml
+    uses: ./.github/workflows/cmake-run-cuda.yml
     with:
       cuda-version: "12.8.0"
 
+  build-hip:
+    needs: build-cuda #  Only attempt HIP after CUDA was succesfull
+    uses: ./.github/workflows/cmake-run-hip.yml
+    with:
+      rocm-version: "6.3.0"
+
+  build-cuda-13-0:
+    if: github.ref_name == 'main'
+    needs: build-cuda
+    uses: ./.github/workflows/cmake-run-cuda.yml
+    with:
+      cuda-version: "13.0.0"
+
   build-cuda-12-6:
+    if: github.ref_name == 'main'
     needs: build-cuda
-    uses: ./.github/workflows/cmake-action.yml
+    uses: ./.github/workflows/cmake-run-cuda.yml
     with:
       cuda-version: "12.6.0"
 
   build-cuda-12-5:
+    if: github.ref_name == 'main'
     needs: build-cuda
-    uses: ./.github/workflows/cmake-action.yml
+    uses: ./.github/workflows/cmake-run-cuda.yml
     with:
       cuda-version: "12.5.0"
+
+  build-hip-6-2:
+    if: github.ref_name == 'main'
+    needs: build-hip
+    uses: ./.github/workflows/cmake-run-hip.yml
+    with:
+      rocm-version: "6.2.0"
+
+  build-hip-6-4:
+    if: github.ref_name == 'main'
+    needs: build-hip
+    uses: ./.github/workflows/cmake-run-hip.yml
+    with:
+      rocm-version: "6.4.0"

.github/workflows/cmake.yml

@@ -107,14 +107,12 @@ KERNEL_FLOAT_BF16_UNARY_FUN(negate, ::__hneg, ::__hneg2)
 #elif KERNEL_FLOAT_IS_HIP
 KERNEL_FLOAT_INLINE __hip_bfloat16 hip_habs(const __hip_bfloat16 a) {
     unsigned short int res = __bfloat16_as_ushort(a);
-    res &= 0x7FFF;
-    return __ushort_as_bfloat16();
+    return __ushort_as_bfloat16(res & 0x7FFF);
 }
 
 KERNEL_FLOAT_INLINE __hip_bfloat16 hip_hneg(const __hip_bfloat16 a) {
     unsigned short int res = __bfloat16_as_ushort(a);
-    res ^= 0x8000;
-    return __ushort_as_bfloat16(res);
+    return __ushort_as_bfloat16(res ^ 0x8000);
 }
 
 KERNEL_FLOAT_INLINE __hip_bfloat162 hip_habs2(const __hip_bfloat162 a) {

include/kernel_float/bf16.h

@@ -19,8 +19,8 @@
     #endif  // __CUDA_ARCH__
 #elif defined(__HIPCC__)
     #define KERNEL_FLOAT_IS_HIP (1)
-    #define KERNEL_FLOAT_DEVICE  __attribute__((always_inline)) __device__
-    #define KERNEL_FLOAT_INLINE  __attribute__((always_inline)) __host__ __device__
+    #define KERNEL_FLOAT_DEVICE  __attribute__((always_inline)) inline __device__
+    #define KERNEL_FLOAT_INLINE  __attribute__((always_inline)) inline __host__ __device__
 
     #ifdef __HIP_DEVICE_COMPILE__
         #define KERNEL_FLOAT_IS_DEVICE (1)

include/kernel_float/macros.h

@@ -100,8 +100,8 @@ struct minmax_tests {
         kf::vec<T, N> a = {x[I]...};
         kf::vec<T, N> b = {y[I]...};
 
-        kf::vec<T, N> lo = min(a, b);
-        kf::vec<T, N> hi = max(a, b);
+        kf::vec<T, N> lo = kernel_float::min(a, b);
+        kf::vec<T, N> hi = kernel_float::max(a, b);
 
         if constexpr (is_one_of<T, double>) {
             ASSERT(equals(fmin(a[I], b[I]), lo[I]) && ...);
@@ -130,7 +130,7 @@ struct cross_test {
     __host__ __device__ void operator()(generator<T> gen) {
         kf::vec<T, 3> a = {T(1.0), T(2.0), T(3.0)};
         kf::vec<T, 3> b = {T(4.0), T(5.0), T(6.0)};
-        kf::vec<T, 3> c = cross(a, b);
+        kf::vec<T, 3> c = kernel_float::cross(a, b);
 
         ASSERT(c[0] == T(-3.0));
         ASSERT(c[1] == T(6.0));

tests/binops.cu

@@ -1,6 +1,8 @@
 #pragma once
 
-#include <cstdint>
+#include <stdint.h>
+#include <math.h>
+#include <tgmath.h>
 
 #include "catch2/catch_all.hpp"
 #include "kernel_float.h"