File size: 3,853 Bytes
2d8a802 | 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95 96 97 98 99 100 101 102 103 104 105 106 107 108 109 110 111 112 113 114 115 116 117 118 119 120 121 122 123 124 125 126 127 128 129 130 131 132 133 134 135 136 137 138 139 140 141 142 143 144 145 146 147 148 149 150 151 152 153 154 | #!/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 |