rocm: add initial exploration code (wip)

Signed-off-by: Daniel Bevenius <[email protected]>

fundamentals/rocm/Makefile

@@ -0,0 +1,5 @@
+
+vector_add: src/vector_add.cpp
+	hipcc -o $@ $<
+
+

fundamentals/rocm/README.md

@@ -0,0 +1,25 @@
+## ROCm exploration
+I currently don't have an AMD GPU so this will only be for learning about the
+libraries and compiling until I can get my hands on one.
+
+## Install ROCm
+
+```console
+wget https://repo.radeon.com/rocm/rocm.gpg.key -O - | \
+    gpg --dearmor | sudo tee /etc/apt/keyrings/rocm.gpg > /dev/null
+echo "deb [arch=amd64 signed-by=/etc/apt/keyrings/rocm.gpg] https://repo.radeon.com/amdgpu/6.1.1/ubuntu jammy main" \
+    | sudo tee /etc/apt/sources.list.d/amdgpu.list
+sudo apt update
+sudo apt install rocm-dev
+```
+Verify that hipcc is installed
+```console
+$ hipcc --version
+HIP version: 6.1.40092-038397aaa
+AMD clang version 17.0.0 (https://github.com/RadeonOpenCompute/llvm-project roc-6.1.1 24154 f53cd7e03908085f4932f7329464cd446426436a)
+Target: x86_64-unknown-linux-gnu
+Thread model: posix
+InstalledDir: /opt/rocm-6.1.1/llvm/bin
+Configuration file: /opt/rocm-6.1.1/lib/llvm/bin/clang++.cfg
+```
+

fundamentals/rocm/set-env.sh

@@ -0,0 +1,3 @@
+#!/bin/bash
+export PATH=/opt/rocm/bin:$PATH
+export LD_LIBRARY_PATH=/opt/rocm/lib:$LD_LIBRARY_PATH

fundamentals/rocm/src/vector_add.cpp

@@ -0,0 +1,37 @@
+#include <hip/hip_runtime.h>
+#include <iostream>
+
+__global__ void vector_add(const float* A, const float* B, float* C, int N) {
+    int idx = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
+    if (idx < N) {
+        C[idx] = A[idx] + B[idx];
+    }
+}
+
+int main() {
+    const int N = 100000;
+    float *A, *B, *C;
+    hipMallocManaged(&A, N * sizeof(float));
+    hipMallocManaged(&B, N * sizeof(float));
+    hipMallocManaged(&C, N * sizeof(float));
+
+    for (int i = 0; i < N; i++) {
+        A[i] = static_cast<float>(i);
+        B[i] = static_cast<float>(i);
+    }
+
+    int blockSize = 256;
+    int numBlocks = (N + blockSize - 1) / blockSize;
+    hipLaunchKernelGGL(vector_add, dim3(numBlocks), dim3(blockSize), 0, 0, A, B, C, N);
+
+    hipDeviceSynchronize();
+
+    std::cout << "C[0] = " << C[0] << std::endl;
+    std::cout << "C[N-1] = " << C[N-1] << std::endl;
+
+    hipFree(A);
+    hipFree(B);
+    hipFree(C);
+
+    return 0;
+}