_dev/debug-build-2-hipcc.sh

#!/usr/bin/env bash

# Debug script 2: HIP compiler basic compilation test

set -euo pipefail

echo "=== HIP Compiler Debug Script 2 ==="
echo "Testing basic HIP compilation capabilities"
echo

# Set ROCm environment variables
export ROCM_PATH="${ROCM_PATH:-/opt/rocm-7.0.1}"
export ROCM_HOME="${ROCM_HOME:-$ROCM_PATH}"
export HIP_PATH="${HIP_PATH:-$ROCM_PATH}"
export HIP_HOME="${HIP_HOME:-$ROCM_PATH}"
export PATH="$ROCM_HOME/bin:$PATH"
export TORCH_HIP_ARCH_LIST="${TORCH_HIP_ARCH_LIST:-gfx942}"
export HSA_OVERRIDE_GFX_VERSION="${HSA_OVERRIDE_GFX_VERSION:-gfx942}"

# Create a simple test directory
mkdir -p /tmp/hip_test
cd /tmp/hip_test

echo "=== Creating Simple HIP Test Program ==="

# Create a minimal HIP program
cat > test_simple.hip << 'EOF'
#include <hip/hip_runtime.h>
#include <iostream>

__global__ void simple_kernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        data[idx] = idx * 2.0f;
    }
}

int main() {
    const int n = 1024;
    const size_t size = n * sizeof(float);

    float* h_data = new float[n];
    float* d_data;

    // Allocate device memory
    hipError_t err = hipMalloc(&d_data, size);
    if (err != hipSuccess) {
        std::cout << "Failed to allocate device memory: " << hipGetErrorString(err) << std::endl;
        return 1;
    }

    // Launch kernel
    dim3 block(256);
    dim3 grid((n + block.x - 1) / block.x);
    hipLaunchKernelGGL(simple_kernel, grid, block, 0, 0, d_data, n);

    // Copy back
    err = hipMemcpy(h_data, d_data, size, hipMemcpyDeviceToHost);
    if (err != hipSuccess) {
        std::cout << "Failed to copy from device: " << hipGetErrorString(err) << std::endl;
        return 1;
    }

    // Check results
    bool success = true;
    for (int i = 0; i < 10; i++) {
        if (h_data[i] != i * 2.0f) {
            success = false;
            break;
        }
    }

    std::cout << "Simple HIP test: " << (success ? "PASSED" : "FAILED") << std::endl;

    hipFree(d_data);
    delete[] h_data;

    return success ? 0 : 1;
}
EOF

echo "=== Testing HIP Compilation ==="
echo "Compiling simple HIP program..."

if hipcc -o test_simple test_simple.hip --amdgpu-target=gfx942 -v; then
    echo "✓ HIP compilation successful"
    echo
    echo "=== Testing HIP Execution ==="
    if ./test_simple; then
        echo "✓ HIP execution successful"
    else
        echo "✗ HIP execution failed"
    fi
else
    echo "✗ HIP compilation failed"
fi

echo
echo "=== Testing HIP with C++ Extensions Flags ==="

# Test with flags similar to what torch uses
echo "Compiling with torch-like flags..."
if hipcc -o test_simple_torch test_simple.hip \
    --amdgpu-target=gfx942 \
    -O3 \
    -std=c++17 \
    -fPIC \
    -shared \
    -v; then
    echo "✓ HIP compilation with torch flags successful"
else
    echo "✗ HIP compilation with torch flags failed"
fi

echo
echo "=== Testing hipblaslt Library ==="
echo "Checking if hipblaslt is available for linking..."

cat > test_hipblaslt.cpp << 'EOF'
#include <iostream>

// Just test if we can link against hipblaslt
int main() {
    std::cout << "hipblaslt linkage test" << std::endl;
    return 0;
}
EOF

if hipcc -o test_hipblaslt test_hipblaslt.cpp -lhipblaslt -v; then
    echo "✓ hipblaslt linkage successful"
    if ./test_hipblaslt; then
        echo "✓ hipblaslt execution successful"
    else
        echo "✗ hipblaslt execution failed"
    fi
else
    echo "✗ hipblaslt linkage failed"
fi

echo
echo "=== Library Search Paths ==="
echo "ROCm lib directory contents:"
ls -la "$ROCM_PATH/lib" | head -20 || echo "Could not list ROCm lib directory"

echo
echo "hipblaslt library files:"
find "$ROCM_PATH" -name "*hipblaslt*" 2>/dev/null || echo "No hipblaslt files found"

echo
echo "=== Debug Script 2 Complete ==="

# Cleanup
cd /
rm -rf /tmp/hip_test