From d6958070cb5ab5c2bb688e29932fc6c2b7ba573b Mon Sep 17 00:00:00 2001 From: Claude Date: Fri, 14 Aug 2026 11:59:51 +0000 Subject: [PATCH] test: probe CUTLASS + __CUDA_ARCH__ on BI-V100 for grouped GEMM feasibility --- qwen3_6_scripts/probe_cutlass_arch.sh | 143 ++++++++++++++++++++++++++ 1 file changed, 143 insertions(+) create mode 100755 qwen3_6_scripts/probe_cutlass_arch.sh diff --git a/qwen3_6_scripts/probe_cutlass_arch.sh b/qwen3_6_scripts/probe_cutlass_arch.sh new file mode 100755 index 00000000..c57b6b18 --- /dev/null +++ b/qwen3_6_scripts/probe_cutlass_arch.sh @@ -0,0 +1,143 @@ +#!/bin/bash +# probe_cutlass_arch.sh — Discover CUTLASS and __CUDA_ARCH__ on BI-V100 +# Run: bash qwen3_6_scripts/probe_cutlass_arch.sh + +echo "=== 1. libcutlass.so symbols (grouped GEMM related) ===" +nm -D /usr/local/corex/lib64/libcutlass.so 2>/dev/null | grep -i "gemm_grouped\|grouped_gemm\|batched_gemm\|GemmGrouped" | head -20 +echo "(total cutlass symbols: $(nm -D /usr/local/corex/lib64/libcutlass.so 2>/dev/null | wc -l))" +echo "" + +echo "=== 2. libcutlass.so ALL exported symbols (first 50) ===" +nm -D /usr/local/corex/lib64/libcutlass.so 2>/dev/null | grep " T " | head -50 +echo "" + +echo "=== 3. CUTLASS headers present? ===" +find /usr/local/corex/include -path "*/cutlass/*" -name "*.h" 2>/dev/null | head -20 +echo "(total cutlass headers: $(find /usr/local/corex/include -path "*/cutlass/*" -name "*.h" 2>/dev/null | wc -l))" +echo "" + +echo "=== 4. CUTLASS grouped GEMM headers? ===" +find /usr/local/corex/include -path "*/cutlass/*" -name "*group*" 2>/dev/null +find /usr/local/corex/include -path "*/cutlass/*" -name "*batch*" 2>/dev/null +echo "" + +echo "=== 5. CUTLASS version ===" +grep -r "CUTLASS_MAJOR\|CUTLASS_MINOR\|CUTLASS_PATCH" /usr/local/corex/include/cutlass/cutlass.h 2>/dev/null | head -5 +echo "" + +echo "=== 6. __CUDA_ARCH__ value on ivcore10 ===" +cat > /tmp/probe_arch.cu << 'CUDA' +#include + +__global__ void print_arch() { +#ifdef __CUDA_ARCH__ + if (threadIdx.x == 0) { + printf("__CUDA_ARCH__ = %d\n", __CUDA_ARCH__); + } +#else + if (threadIdx.x == 0) { + printf("__CUDA_ARCH__ not defined (host code)\n"); + } +#endif +} + +// Host-side macro check +int main() { + printf("Host: __CUDA_ARCH__ "); +#ifdef __CUDA_ARCH__ + printf("= %d\n", __CUDA_ARCH__); +#else + printf("not defined (expected for host)\n"); +#endif + + print_arch<<<1, 1>>>(); + cudaDeviceSynchronize(); + + // Also check compute capability via runtime API + cudaDeviceProp prop; + cudaGetDeviceProperties(&prop, 0); + printf("cudaDeviceProp: major=%d minor=%d\n", prop.major, prop.minor); + printf("Device name: %s\n", prop.name); + printf("SM count: %d\n", prop.multiProcessorCount); + printf("Shared mem per block: %zu\n", prop.sharedMemPerBlock); + printf("Shared mem per SM: %zu\n", prop.sharedMemPerMultiprocessor); + printf("Warp size: %d\n", prop.warpSize); + printf("Max threads per block: %d\n", prop.maxThreadsPerBlock); + printf("Clock rate: %d kHz\n", prop.clockRate); + printf("Memory clock: %d kHz\n", prop.memoryClockRate); + printf("Memory bus width: %d bits\n", prop.memoryBusWidth); + printf("L2 cache: %d bytes\n", prop.l2CacheSize); + printf("Total global mem: %zu bytes\n", prop.totalGlobalMem); + + return 0; +} +CUDA + +echo "Compiling..." +/usr/local/corex/bin/clang++ --cuda-gpu-arch=ivcore10 --cuda-path=/usr/local/corex \ + -I/usr/local/corex/include -L/usr/local/corex/lib64 -lcudart \ + /tmp/probe_arch.cu -o /tmp/probe_arch 2>&1 + +if [ -f /tmp/probe_arch ]; then + echo "Compile: SUCCESS" + /tmp/probe_arch +else + echo "Compile: FAILED" +fi +echo "" + +echo "=== 7. PyTorch CUDA arch ===" +python3 -c " +import torch +print(f'torch.cuda.get_device_capability() = {torch.cuda.get_device_capability()}') +print(f'torch.cuda.get_device_name() = {torch.cuda.get_device_name()}') +print(f'torch.cuda.get_arch_list() = {torch.cuda.get_arch_list()}') +print(f'torch.version.cuda = {torch.version.cuda}') +" 2>/dev/null +echo "" + +echo "=== 8. cublasLt symbols (batched GEMM) ===" +nm -D /usr/local/corex/lib64/libcublasLt.so 2>/dev/null | grep -i "matmul\|batch\|group" | head -20 +echo "(total cublasLt symbols: $(nm -D /usr/local/corex/lib64/libcublasLt.so 2>/dev/null | wc -l))" +echo "" + +echo "=== 9. libcuinfer.so (if it has grouped GEMM) ===" +nm -D /usr/local/corex/lib64/libcuinfer.so 2>/dev/null | grep -i "gemm\|matmul\|expert\|moe" | head -20 +echo "" + +echo "=== 10. ixformer _ixformer_torch.so — check for WMMA/gemm symbols ===" +TORCH_SO=$(find /usr/local/corex -name "_ixformer_torch*.so" 2>/dev/null | head -1) +if [ -n "$TORCH_SO" ]; then + nm -D "$TORCH_SO" 2>/dev/null | grep -i "wmma\|gemm\|moe\|expert\|group" | head -30 +fi +echo "" + +echo "=== 11. CUTLASS GEMM templates check ===" +ls /usr/local/corex/include/cutlass/gemm/ 2>/dev/null +ls /usr/local/corex/include/cutlass/gemm/device/ 2>/dev/null +echo "" + +echo "=== 12. Quick CUTLASS compile test ===" +cat > /tmp/test_cutlass.cu << 'CUDA' +#include +#include + +int main() { + printf("CUTLASS version: %d.%d.%d\n", + CUTLASS_MAJOR, CUTLASS_MINOR, CUTLASS_PATCH); + return 0; +} +CUDA + +/usr/local/corex/bin/clang++ --cuda-gpu-arch=ivcore10 --cuda-path=/usr/local/corex \ + -I/usr/local/corex/include -L/usr/local/corex/lib64 -lcudart -lcutlass \ + /tmp/test_cutlass.cu -o /tmp/test_cutlass 2>&1 + +if [ -f /tmp/test_cutlass ]; then + echo "CUTLASS compile: SUCCESS" + /tmp/test_cutlass +else + echo "CUTLASS compile: FAILED (may need different include path)" + # Try alternative paths + find /usr/local/corex -name "cutlass.h" -path "*/cutlass/*" 2>/dev/null +fi