Static Device Library support + llvm configure fix

scripts/configure_llvm.sh

@@ -188,31 +188,32 @@ else
 fi
 
 # Add build type condition
+# Forcing the use of gcc and g++ to avoid issues with intel compilers
+COMMON_CMAKE_OPTIONS=(
+  -DCMAKE_CXX_COMPILER=g++ \
+  -DCMAKE_C_COMPILER=gcc \
+  -DCMAKE_INSTALL_PREFIX=${INSTALL_DIR} \
+  -DCMAKE_BUILD_TYPE=Release \
+  -DLLVM_ENABLE_PROJECTS="clang;openmp;clang-tools-extra" \
+  -DLLVM_TARGETS_TO_BUILD=host \
+  -DLLVM_EXPERIMENTAL_TARGETS_TO_BUILD="SPIRV" \
+  -DLLVM_ENABLE_ASSERTIONS=On \
+  -DLLVM_BINUTILS_INCDIR=${BINUTILS_HEADER_DIR}
+  -DCMAKE_CXX_LINK_FLAGS="-Wl,-rpath,${gcc_base_path}/lib64 -L${gcc_base_path}/lib64"
+)
+
 if [ "$LINK_TYPE" == "static" ]; then
   cmake ../  \
-    -DCMAKE_INSTALL_PREFIX=${INSTALL_DIR} \
-    -DCMAKE_BUILD_TYPE=Release \
-    -DLLVM_ENABLE_PROJECTS="clang;openmp;clang-tools-extra" \
-    -DLLVM_TARGETS_TO_BUILD=host \
-    -DLLVM_EXPERIMENTAL_TARGETS_TO_BUILD="SPIRV" \
-    -DLLVM_ENABLE_ASSERTIONS=On \
-    -DLLVM_BINUTILS_INCDIR=${BINUTILS_HEADER_DIR} \
-    -DCMAKE_CXX_LINK_FLAGS="-Wl,-rpath,${gcc_base_path}/lib64 -L${gcc_base_path}/lib64"
+    "${COMMON_CMAKE_OPTIONS[@]}" 
 elif [ "$LINK_TYPE" == "dynamic" ]; then
   cmake ../ \
-    -DCMAKE_INSTALL_PREFIX=${INSTALL_DIR} \
+    "${COMMON_CMAKE_OPTIONS[@]}" \
     -DCMAKE_INSTALL_RPATH=${INSTALL_DIR}/lib \
-    -DLLVM_ENABLE_PROJECTS="clang;openmp;clang-tools-extra" \
-    -DLLVM_TARGETS_TO_BUILD=host \
-    -DLLVM_EXPERIMENTAL_TARGETS_TO_BUILD="SPIRV" \
     -DLLVM_LINK_LLVM_DYLIB=ON \
-    -DLLVM_BUILD_LLVM_DYLIB=ON \
-    -DLLVM_PARALLEL_LINK_JOBS=2 \
-    -DCMAKE_BUILD_TYPE=Release \
-    -DCMAKE_CXX_LINK_FLAGS="-Wl,-rpath,${gcc_base_path}/lib64 -L${gcc_base_path}/lib64" \
-    -DLLVM_BINUTILS_INCDIR=${BINUTILS_HEADER_DIR} \
-    -DLLVM_ENABLE_ASSERTIONS=On
+    -DLLVM_BUILD_LLVM_DYLIB=ON
 else
   echo "Invalid link_type. Must be 'static' or 'dynamic'."
   exit 1
 fi
+
+# Make sure ninja is in the path

tests/compiler/CMakeLists.txt

@@ -143,6 +143,7 @@ add_shell_test(TestHipcc692Regression.bash)
 add_shell_test(TestHipccFileOrderPreservation.bash)
 add_shell_test(TestHipcc945MacroSpace.bash)
 add_shell_test(TestHipccArgOrder.bash)
+add_subdirectory(rdcLink)
 add_test(NAME "TestHipccMultiSource" COMMAND 
   ${CMAKE_BINARY_DIR}/bin/hipcc ${CMAKE_CURRENT_SOURCE_DIR}/TestHipccCompileThenLinkMain.cpp ${CMAKE_CURRENT_SOURCE_DIR}/TestHipccCompileThenLinkKernel.cpp -o TestHipccMultiSource)
 

tests/compiler/rdcLink/CMakeLists.txt

@@ -0,0 +1 @@
+add_shell_test(TestStaticLibRDC.bash)

tests/compiler/rdcLink/TestStaticLibRDC.bash

@@ -0,0 +1,49 @@
+#!/bin/bash
+
+# Test description:
+# This test verifies that building a static library with device code using
+# -fgpu-rdc fails as described in https://github.com/CHIP-SPV/chipStar/issues/984
+
+# Exit script on error
+set -eu
+
+# CMake substituted variables
+SRC_DIR="@CMAKE_CURRENT_SOURCE_DIR@"
+OUT_DIR="@CMAKE_CURRENT_BINARY_DIR@/@TEST_NAME@.d"
+HIPCC="@CMAKE_BINARY_DIR@/bin/hipcc"
+TEST_NAME="@TEST_NAME@" # Make sure TEST_NAME is available if needed later
+
+# Create output directory
+mkdir -p "${OUT_DIR}"
+
+# Compile the device code files
+${HIPCC} -fgpu-rdc -fPIC -I"${SRC_DIR}" -c "${SRC_DIR}/k.cu" -o "${OUT_DIR}/k.o"
+${HIPCC} -fgpu-rdc -fPIC -I"${SRC_DIR}" -c "${SRC_DIR}/k1.cu" -o "${OUT_DIR}/k1.o"
+
+# Create the static library
+ar rcs "${OUT_DIR}/libk.a" "${OUT_DIR}/k.o" "${OUT_DIR}/k1.o"
+
+# Compile the main host file
+${HIPCC} -fgpu-rdc -I"${SRC_DIR}" -c "${SRC_DIR}/t.cpp" -o "${OUT_DIR}/t.o"
+
+# Link the main file and the static library
+${HIPCC} --save-temps -v -fgpu-rdc --hip-link "${OUT_DIR}/t.o" "${OUT_DIR}/libk.a" -o "${OUT_DIR}/TestStaticLibRDC" 
+echo "TestStaticLibRDC.log: ${OUT_DIR}/TestStaticLibRDC.log"
+${OUT_DIR}/TestStaticLibRDC
+
+# Attempt to run the executable - we expect this to fail with a specific error
+# Use run_and_check_error.bash helper if available, otherwise grep for the error
+RUN_EXEC="${OUT_DIR}/TestStaticLibRDC"
+
+# Run the executable, capture stderr, and check for success and absence of errors
+echo "Running: ${RUN_EXEC}"
+STDERR_OUTPUT=$("${RUN_EXEC}" 2>&1) || true # Run and capture stderr, allow non-zero exit code here
+
+
+# Check for explicit error messages in stderr
+if echo "${STDERR_OUTPUT}" | grep -qE 'CHIP error|hipError'; then
+  echo "Test FAILED: Error messages found in output."
+  echo "Output:"
+  echo "${STDERR_OUTPUT}"
+  exit 1
+fi

tests/compiler/rdcLink/k.cu

@@ -0,0 +1,20 @@
+#include "k.h"
+
+#define HIP_ASSERT(x) (assert((x)==hipSuccess))
+
+__global__ void saxpy2(int n, float a, float *x, float *y)
+{
+  int i = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
+  if (i < n) y[i] = a*x[i] + y[i];
+}
+
+__device__ int device_square(int x)
+{
+    return x * x;
+}
+
+void test(float *d_x,float *d_y, float *x, float *y, int N )
+{
+  hipLaunchKernelGGL(saxpy2,(N+255)/256, 256,0,0,N, 2.0f, d_x, d_y );
+  hipDeviceSynchronize();
+} 

tests/compiler/rdcLink/k.h

@@ -0,0 +1,11 @@
+#ifndef TEST_STATIC_LIB_RDC_K_H
+#define TEST_STATIC_LIB_RDC_K_H
+
+#include "hip/hip_runtime.h"
+#include <cassert>
+
+__device__ int device_square(int x);
+void test(float *d_x,float *d_y, float *x, float *y, int N );
+void test2(float *d_x,float *d_y, float *x, float *y, int N );
+
+#endif // TEST_STATIC_LIB_RDC_K_H 

tests/compiler/rdcLink/k1.cu

@@ -0,0 +1,25 @@
+#include "k.h"
+
+#define HIP_ASSERT(x) (assert((x)==hipSuccess))
+
+__global__ void saxpy3(int n, float a, float *x, float *y)
+{
+  int i = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
+  if (i < n) y[i] = a*x[i] + y[i];
+}
+
+__global__ void square_kernel(int* out, const int* in, const unsigned int size)
+{
+    const unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
+    if(tid < size)
+    {
+        out[tid] = device_square(in[tid]);
+    }
+}
+
+void test2(float *d_x,float *d_y, float *x, float *y, int N )
+{
+  hipLaunchKernelGGL(saxpy3,(N+255)/256, 256,0,0,N, 2.0f, d_x, d_y );
+  hipDeviceSynchronize();
+  float pABC[] = {1, 2, 3};
+} 

tests/compiler/rdcLink/t.cpp

@@ -0,0 +1,33 @@
+#include <stdio.h>
+#include <math.h>
+#include <stdlib.h> // Added for malloc
+#include "k.h"
+
+#define HIP_ASSERT(x) (assert((x)==hipSuccess))
+
+int main(void)
+{
+  int N = 1<<20;
+  float *x, *y, *d_x, *d_y;
+
+  x = (float*)malloc(N*sizeof(float));
+  y = (float*)malloc(N*sizeof(float));
+
+  HIP_ASSERT(hipMalloc(&d_x, N*sizeof(float)));
+  HIP_ASSERT(hipMalloc(&d_y, N*sizeof(float)));
+
+  for (int i = 0; i < N; i++) {
+    x[i] = 1.0f;
+    y[i] = 2.0f;
+  }
+
+  HIP_ASSERT(hipMemcpy(d_x, x, N*sizeof(float), hipMemcpyHostToDevice));
+  HIP_ASSERT(hipMemcpy(d_y, y, N*sizeof(float), hipMemcpyHostToDevice));
+
+  test( d_x, d_y, x, y, N);
+  test2( d_x, d_y, x, y, N);
+
+  // TODO: Add hipFree and free calls for proper cleanup
+
+  return 0; // Added return statement
+}