Run HIP integration tests

Signed-off-by: Gavin Zhao <git@gzgz.dev>

.github/workflows/integration-tests.yml

@@ -472,6 +472,78 @@ jobs:
 
           ${SCCACHE_PATH} --show-stats | grep -e "Cache hits\s*[1-9]"
 
+  hip:
+    # Probably wouldn't matter anyway since we run in a container, but staying 
+    # close to the version is better than not.
+    runs-on: ubuntu-22.04
+    needs: build
+    container:
+      image: rocm/dev-ubuntu-22.04:6.0
+
+    env:
+      # SCCACHE_GHA_ENABLED: "on"
+      ROCM_PATH: "/usr"
+
+    steps:
+      - uses: actions/checkout@v4
+
+      # I don't want to break the cache during testing. Will turn on after I
+      # make sure it's working.
+      #
+      # - name: Configure Cache Env
+      #   uses: actions/github-script@v7
+      #   with:
+      #     script: |
+      #       core.exportVariable('ACTIONS_CACHE_URL', process.env.ACTIONS_CACHE_URL || '');
+      #       core.exportVariable('ACTIONS_RUNTIME_TOKEN', process.env.ACTIONS_RUNTIME_TOKEN || '')
+
+      # - name: Configure ROCm Env
+      #   uses: actions/github-script@v7
+      #   with:
+      #     script: |
+      #       core.exportVariable('ROCM_PATH', process.env.ROCM_PATH || '');
+
+      - uses: actions/download-artifact@v4
+        with:
+          name: integration-tests
+          path: /home/runner/.cargo/bin/
+      - name: Chmod for binary
+        run: chmod +x ${SCCACHE_PATH}
+
+      - name: Install dependencies
+        shell: bash
+        run: |
+          ## Install dependencies
+          sudo apt-get update
+          sudo apt-get install -y cmake
+
+      # Ensure that HIPCC isn't already borken
+      - name: Sanity Check
+        run: |
+          hipcc -o vectoradd_hip --offload-arch=gfx900 tests/cmake-hip/vectoradd_hip.cpp
+
+      - name: Test
+        run: |
+          cmake -B build -S tests/cmake-hip -DCMAKE_HIP_COMPILER_LAUNCHER=${SCCACHE_PATH} -DCMAKE_HIP_ARCHITECTURES=gfx900
+          cmake --build build
+
+      - name: Output
+        run: |
+          ${SCCACHE_PATH} --show-stats
+
+      - name: Test Twice for Cache Read
+        run: |
+          rm -rf build
+          cmake -B build -S tests/cmake-hip -DCMAKE_HIP_COMPILER_LAUNCHER=${SCCACHE_PATH} -DCMAKE_HIP_ARCHITECTURES=gfx900
+          cmake --build build
+
+      - name: Output
+        run: |
+          ${SCCACHE_PATH} --show-stats
+
+          ${SCCACHE_PATH} --show-stats | grep -e "Cache hits\s*[1-9]"
+
+
   gcc:
     runs-on: ubuntu-latest
     needs: build

tests/cmake-hip/CMakeLists.txt

@@ -0,0 +1,6 @@
+cmake_minimum_required(VERSION 3.10)
+
+project(myproject LANGUAGES CXX HIP)
+
+add_library(vectoradd_hip vectoradd_hip.cpp)
+set_source_files_properties(vectoradd_hip.cpp PROPERTIES LANGUAGE HIP)

tests/cmake-hip/vectoradd_hip.cpp

@@ -0,0 +1,150 @@
+/*
+Copyright (c) 2015-2016 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 <assert.h>
+#include <stdio.h>
+#include <algorithm>
+#include <stdlib.h>
+#include<iostream>
+#include "hip/hip_runtime.h"
+
+
+#define HIP_ASSERT(x) (assert((x)==hipSuccess))
+
+
+#define WIDTH     1024
+#define HEIGHT    1024
+
+#define NUM       (WIDTH*HEIGHT)
+
+#define THREADS_PER_BLOCK_X  16
+#define THREADS_PER_BLOCK_Y  16
+#define THREADS_PER_BLOCK_Z  1
+
+__global__ void 
+vectoradd_float(float* __restrict__ a, const float* __restrict__ b, const float* __restrict__ c, int width, int height) 
+
+  {
+
+      int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
+      int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
+
+      int i = y * width + x;
+      if ( i < (width * height)) {
+        a[i] = b[i] + c[i];
+      }
+
+
+
+  }
+
+#if 0
+__kernel__ void vectoradd_float(float* a, const float* b, const float* c, int width, int height) {
+
+
+  int x = blockDimX * blockIdx.x + threadIdx.x;
+  int y = blockDimY * blockIdy.y + threadIdx.y;
+
+  int i = y * width + x;
+  if ( i < (width * height)) {
+    a[i] = b[i] + c[i];
+  }
+}
+#endif
+
+using namespace std;
+
+int main() {
+
+  float* hostA;
+  float* hostB;
+  float* hostC;
+
+  float* deviceA;
+  float* deviceB;
+  float* deviceC;
+
+  hipDeviceProp_t devProp;
+  hipGetDeviceProperties(&devProp, 0);
+  cout << " System minor " << devProp.minor << endl;
+  cout << " System major " << devProp.major << endl;
+  cout << " agent prop name " << devProp.name << endl;
+
+
+
+  cout << "hip Device prop succeeded " << endl ;
+
+
+  int i;
+  int errors;
+
+  hostA = (float*)malloc(NUM * sizeof(float));
+  hostB = (float*)malloc(NUM * sizeof(float));
+  hostC = (float*)malloc(NUM * sizeof(float));
+
+  // initialize the input data
+  for (i = 0; i < NUM; i++) {
+    hostB[i] = (float)i;
+    hostC[i] = (float)i*100.0f;
+  }
+
+  HIP_ASSERT(hipMalloc((void**)&deviceA, NUM * sizeof(float)));
+  HIP_ASSERT(hipMalloc((void**)&deviceB, NUM * sizeof(float)));
+  HIP_ASSERT(hipMalloc((void**)&deviceC, NUM * sizeof(float)));
+
+  HIP_ASSERT(hipMemcpy(deviceB, hostB, NUM*sizeof(float), hipMemcpyHostToDevice));
+  HIP_ASSERT(hipMemcpy(deviceC, hostC, NUM*sizeof(float), hipMemcpyHostToDevice));
+
+
+  hipLaunchKernelGGL(vectoradd_float, 
+                  dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
+                  dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
+                  0, 0,
+                  deviceA ,deviceB ,deviceC ,WIDTH ,HEIGHT);
+
+
+  HIP_ASSERT(hipMemcpy(hostA, deviceA, NUM*sizeof(float), hipMemcpyDeviceToHost));
+
+  // verify the results
+  errors = 0;
+  for (i = 0; i < NUM; i++) {
+    if (hostA[i] != (hostB[i] + hostC[i])) {
+      errors++;
+    }
+  }
+  if (errors!=0) {
+    printf("FAILED: %d errors\n",errors);
+  } else {
+      printf ("PASSED!\n");
+  }
+
+  HIP_ASSERT(hipFree(deviceA));
+  HIP_ASSERT(hipFree(deviceB));
+  HIP_ASSERT(hipFree(deviceC));
+
+  free(hostA);
+  free(hostB);
+  free(hostC);
+
+  //hipResetDefaultAccelerator();
+
+  return errors;
+}