CodeQL / Analyze (actions) (push) Canceled after 0s
CodeQL / Analyze (python) (push) Canceled after 0s
pre-commit / pre-commit (push) Canceled after 0s
Multi-Arch CI / setup (push) Canceled after 0s
Publish build_manylinux_x86_64 images / publish_build_manylinux_x86_64 (push) Canceled after 0s
Setup (Multi-Arch) / setup (push) Canceled after 0s
Multi-Arch CI / Linux::${{ fromJSON(needs.setup.outputs.linux_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubi10 image / publish_no_rocm_image_ubi10 (push) Canceled after 0s
Multi-Arch CI / Windows::${{ fromJSON(needs.setup.outputs.windows_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Multi-Arch CI / Manifest Diff (push) Canceled after 0s
TheRock Manifest Diff Report / Generate Manifest Diff Report (push) Canceled after 0s
Multi-Arch CI / CI Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_ocl_rt images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_rocgdb images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Unit Tests / Unit Tests :: ubuntu-24.04 (push) Canceled after 0s
Unit Tests / Unit Tests :: windows-2022 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_media (push) Canceled after 0s
Unit Tests / Unit Tests Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_openmpi (push) Canceled after 0s
300 lines
9.5 KiB
Python
300 lines
9.5 KiB
Python
"""Tests for CCOB parser module."""
|
|
|
|
import struct
|
|
import subprocess
|
|
from pathlib import Path
|
|
|
|
import pytest
|
|
import zstandard as zstd
|
|
|
|
from rocm_kpack.ccob_parser import (
|
|
CCOBHeader,
|
|
UncompressedBundle,
|
|
decompress_ccob,
|
|
extract_code_objects_from_fatbin,
|
|
list_ccob_targets,
|
|
parse_ccob_file,
|
|
parse_fatbin_data,
|
|
)
|
|
|
|
|
|
def test_ccob_header_parse():
|
|
"""Test CCOB header parsing."""
|
|
# Construct a minimal CCOB v3 header
|
|
header = (
|
|
b"CCOB" # Magic
|
|
b"\x03\x00\x01\x00" # Version 3, compression method 1 (zstd)
|
|
b"\x0c\x8b\x12\x00\x00\x00\x00\x00" # totalSize: 1,215,244
|
|
b"\x38\x7f\xc6\x00\x00\x00\x00\x00" # uncompressedSize: 13,008,696
|
|
b"\xde\x30\xdc\xec\xea\x03\x74\xc3" # hash
|
|
)
|
|
|
|
parsed = CCOBHeader.parse(header)
|
|
assert parsed.magic == "CCOB"
|
|
assert parsed.version == 3
|
|
assert parsed.compression_method == 1
|
|
assert parsed.total_size == 1_215_244
|
|
assert parsed.uncompressed_size == 13_008_696
|
|
|
|
|
|
def test_ccob_header_invalid_magic():
|
|
"""Test that invalid magic raises ValueError."""
|
|
bad_header = b"XXXX" + b"\x00" * 28
|
|
|
|
with pytest.raises(ValueError, match="Invalid magic"):
|
|
CCOBHeader.parse(bad_header)
|
|
|
|
|
|
def test_ccob_header_too_short():
|
|
"""Test that short header raises ValueError."""
|
|
short_header = b"CCOB" + b"\x00" * 20
|
|
|
|
with pytest.raises(ValueError, match="Header too short"):
|
|
CCOBHeader.parse(short_header)
|
|
|
|
|
|
def test_decompress_ccob_with_real_binary(tmp_path: Path, toolchain):
|
|
"""Test CCOB decompression with a real ROCm binary."""
|
|
# Find a real ROCm library with CCOB bundle
|
|
rocm_lib = Path("/home/stella/workspace/rocm/gfx1100/lib/librocblas.so.5")
|
|
|
|
if not rocm_lib.exists():
|
|
pytest.skip("librocblas.so.5 not found")
|
|
|
|
# Extract .hip_fatbin section
|
|
fatbin_path = tmp_path / "librocblas_fatbin.o"
|
|
result = subprocess.run(
|
|
["objcopy", "--dump-section", f".hip_fatbin={fatbin_path}", str(rocm_lib)],
|
|
capture_output=True,
|
|
)
|
|
|
|
if result.returncode != 0:
|
|
pytest.skip("Could not extract .hip_fatbin section")
|
|
|
|
# Parse and decompress
|
|
bundle = parse_ccob_file(fatbin_path)
|
|
|
|
# Verify structure
|
|
assert bundle.magic.startswith("__CLANG_OFFLOAD_BUNDLE__")
|
|
assert bundle.num_entries > 0
|
|
|
|
# Should have at least one device target
|
|
triples = bundle.list_triples()
|
|
device_targets = [t for t in triples if "hipv4-amdgcn" in t]
|
|
assert len(device_targets) > 0
|
|
|
|
# Verify we can extract code objects
|
|
for triple in device_targets:
|
|
code_obj = bundle.get_code_object(triple)
|
|
assert code_obj is not None
|
|
assert len(code_obj) > 0
|
|
|
|
|
|
def test_list_ccob_targets_with_real_binary(tmp_path: Path):
|
|
"""Test listing CCOB targets."""
|
|
rocm_lib = Path("/home/stella/workspace/rocm/gfx1100/lib/librocblas.so.5")
|
|
|
|
if not rocm_lib.exists():
|
|
pytest.skip("librocblas.so.5 not found")
|
|
|
|
# Extract .hip_fatbin section
|
|
fatbin_path = tmp_path / "librocblas_fatbin.o"
|
|
result = subprocess.run(
|
|
["objcopy", "--dump-section", f".hip_fatbin={fatbin_path}", str(rocm_lib)],
|
|
capture_output=True,
|
|
)
|
|
|
|
if result.returncode != 0:
|
|
pytest.skip("Could not extract .hip_fatbin section")
|
|
|
|
# List targets
|
|
data = fatbin_path.read_bytes()
|
|
targets = list_ccob_targets(data)
|
|
|
|
# Should have host + device targets
|
|
assert len(targets) > 0
|
|
assert any("host-" in t for t in targets)
|
|
assert any("hipv4-amdgcn" in t for t in targets)
|
|
|
|
|
|
def test_decompress_respects_total_size():
|
|
"""Test that decompression respects totalSize field, not buffer size.
|
|
|
|
This is the key fix that clang-offload-bundler is missing.
|
|
"""
|
|
# This test would need a crafted CCOB with padding, which we have
|
|
# in librocblas.so.5. The test above verifies this works.
|
|
pass
|
|
|
|
|
|
def test_uncompressed_bundle_get_code_object():
|
|
"""Test getting specific code object from bundle."""
|
|
# Create a minimal uncompressed bundle structure
|
|
magic = b"__CLANG_OFFLOAD_BUNDLE__\x00"[:24]
|
|
num_entries = (1).to_bytes(8, "little")
|
|
|
|
# Single entry: offset=100, size=10, triple="test-triple"
|
|
triple = b"test-triple"
|
|
entry = (
|
|
(100).to_bytes(8, "little") # offset
|
|
+ (10).to_bytes(8, "little") # size
|
|
+ len(triple).to_bytes(8, "little") # triple_size
|
|
+ triple
|
|
)
|
|
|
|
# Padding to offset 100, then code object data
|
|
padding = b"\x00" * (100 - 32 - len(entry))
|
|
code_data = b"0123456789"
|
|
|
|
bundle_data = magic + num_entries + entry + padding + code_data
|
|
|
|
# Parse
|
|
bundle = UncompressedBundle.parse(bundle_data)
|
|
|
|
assert bundle.num_entries == 1
|
|
assert bundle.entries[0].triple == "test-triple"
|
|
|
|
# Get code object
|
|
code_obj = bundle.get_code_object("test-triple")
|
|
assert code_obj == b"0123456789"
|
|
|
|
# Try non-existent triple
|
|
assert bundle.get_code_object("nonexistent") is None
|
|
|
|
|
|
# ---------------------------------------------------------------------------
|
|
# Helpers for building synthetic CCOB bundles
|
|
# ---------------------------------------------------------------------------
|
|
|
|
|
|
def _make_uncompressed_bundle(entries: list[tuple[str, bytes]]) -> bytes:
|
|
"""Build an uncompressed offload bundle from (triple, data) pairs."""
|
|
magic = b"__CLANG_OFFLOAD_BUNDLE__"
|
|
num_entries = struct.pack("<Q", len(entries))
|
|
|
|
# Compute header size to know where code objects start
|
|
header_size = 24 + 8 # magic + num_entries
|
|
for triple, _ in entries:
|
|
header_size += 8 + 8 + 8 + len(triple) # offset + size + triple_len + triple
|
|
|
|
# Build entry descriptors and concatenate code object data
|
|
descriptors = b""
|
|
code_blobs = b""
|
|
for triple, data in entries:
|
|
offset = header_size + len(code_blobs)
|
|
descriptors += struct.pack("<QQQ", offset, len(data), len(triple))
|
|
descriptors += triple.encode("ascii")
|
|
code_blobs += data
|
|
|
|
return magic + num_entries + descriptors + code_blobs
|
|
|
|
|
|
def _make_ccob(bundle_data: bytes) -> bytes:
|
|
"""Wrap an uncompressed bundle in a CCOB compressed envelope."""
|
|
cctx = zstd.ZstdCompressor()
|
|
compressed = cctx.compress(bundle_data)
|
|
|
|
# CCOB v3 header: 4B magic + 2B version + 2B method + 8B totalSize + 8B uncompSize + 8B hash
|
|
header_size = 32
|
|
total_size = header_size + len(compressed)
|
|
header = (
|
|
b"CCOB"
|
|
+ struct.pack("<HH", 3, 1) # version=3, method=1 (zstd)
|
|
+ struct.pack("<Q", total_size)
|
|
+ struct.pack("<Q", len(bundle_data))
|
|
+ struct.pack("<Q", 0) # hash (unused for decompression)
|
|
)
|
|
return header + compressed
|
|
|
|
|
|
# ---------------------------------------------------------------------------
|
|
# Multi-CCOB tests
|
|
# ---------------------------------------------------------------------------
|
|
|
|
|
|
def test_parse_fatbin_data_single_ccob():
|
|
"""Single CCOB round-trips correctly."""
|
|
bundle_bytes = _make_uncompressed_bundle(
|
|
[
|
|
("host-x86_64-unknown-linux-gnu-", b""),
|
|
("hipv4-amdgcn-amd-amdhsa--gfx942", b"CODE_OBJ_A"),
|
|
]
|
|
)
|
|
ccob = _make_ccob(bundle_bytes)
|
|
|
|
bundles = parse_fatbin_data(ccob)
|
|
assert len(bundles) == 1
|
|
assert bundles[0].num_entries == 2
|
|
|
|
|
|
def test_parse_fatbin_data_multiple_ccobs():
|
|
"""Multiple back-to-back CCOBs are all extracted.
|
|
|
|
Matches the layout produced by the linker for multi-TU HIP binaries.
|
|
Verifies the fix for the bug where only the first CCOB was parsed.
|
|
See: llvm/lib/Object/OffloadBundle.cpp:33-99 (extractOffloadBundle)
|
|
"""
|
|
bundle_a = _make_uncompressed_bundle(
|
|
[
|
|
("host-x86_64-unknown-linux-gnu-", b""),
|
|
("hipv4-amdgcn-amd-amdhsa--gfx942", b"KERNEL_A"),
|
|
]
|
|
)
|
|
bundle_b = _make_uncompressed_bundle(
|
|
[
|
|
("host-x86_64-unknown-linux-gnu-", b""),
|
|
("hipv4-amdgcn-amd-amdhsa--gfx942", b"KERNEL_B"),
|
|
]
|
|
)
|
|
bundle_c = _make_uncompressed_bundle(
|
|
[
|
|
("host-x86_64-unknown-linux-gnu-", b""),
|
|
("hipv4-amdgcn-amd-amdhsa--gfx942", b"KERNEL_C"),
|
|
]
|
|
)
|
|
|
|
fatbin = _make_ccob(bundle_a) + _make_ccob(bundle_b) + _make_ccob(bundle_c)
|
|
|
|
bundles = parse_fatbin_data(fatbin)
|
|
assert len(bundles) == 3
|
|
|
|
# Verify we can extract each code object
|
|
objs = extract_code_objects_from_fatbin(fatbin)
|
|
assert len(objs) == 3
|
|
assert objs[0].data == b"KERNEL_A"
|
|
assert objs[1].data == b"KERNEL_B"
|
|
assert objs[2].data == b"KERNEL_C"
|
|
|
|
|
|
def test_parse_fatbin_data_ccobs_with_padding():
|
|
"""CCOBs separated by page-alignment padding (zeros).
|
|
|
|
Real linker output page-aligns each CCOB. The parser must skip padding
|
|
bytes between entries, matching LLVM's magic-scanning approach.
|
|
"""
|
|
ccob_a = _make_ccob(
|
|
_make_uncompressed_bundle([("hipv4-amdgcn-amd-amdhsa--gfx942", b"A" * 100)])
|
|
)
|
|
ccob_b = _make_ccob(
|
|
_make_uncompressed_bundle([("hipv4-amdgcn-amd-amdhsa--gfx942", b"B" * 100)])
|
|
)
|
|
|
|
# Pad first CCOB to next 4096-byte boundary
|
|
padded_size = ((len(ccob_a) + 4095) // 4096) * 4096
|
|
padding = b"\x00" * (padded_size - len(ccob_a))
|
|
fatbin = ccob_a + padding + ccob_b
|
|
|
|
objs = extract_code_objects_from_fatbin(fatbin)
|
|
assert len(objs) == 2
|
|
assert objs[0].data == b"A" * 100
|
|
assert objs[1].data == b"B" * 100
|
|
|
|
|
|
def test_parse_fatbin_data_empty_raises():
|
|
"""Empty or unrecognized data raises ValueError."""
|
|
with pytest.raises(ValueError, match="Unrecognized fatbin format"):
|
|
parse_fatbin_data(b"")
|
|
|
|
with pytest.raises(ValueError, match="Unrecognized fatbin format"):
|
|
parse_fatbin_data(b"\x00" * 100)
|