#! /usr/bin/make -f
# Use ROCm's own compiler (package clang-rocm) rather than the distro's
# generic llvm-* package: the bulk of this build is actual GPU device code
# (Tensile-generated kernels, device-library extops, matrix-transform), and
# ROCm's LLVM fork carries AMDGPU-specific ISA/assembler extensions and
# device library layout that generic llvm-* does not have (or is not
# compatible with) -- confirmed by chasing this per-callsite across HIP
# device compiles, raw hsaco builds, GPU assembly, and codegen-script
# toolchain args before concluding the shared root cause. Absolute paths
# sidestep PATH-ordering concerns entirely (build-time PATH is not
# guaranteed to carry over to autopkgtest time).
export CC=/usr/lib/rocm/llvm/bin/amdclang
export CXX=/usr/lib/rocm/llvm/bin/amdclang++
export DEB_BUILD_MAINT_OPTIONS = hardening=+all optimize=-lto
export DEB_CXXFLAGS_MAINT_PREPEND = -gz
export VERBOSE=1
export DPKG_GENSYMBOLS_CHECK_LEVEL = 4
export PATH:=/usr/lib/rocm/llvm/bin:$(PATH)
#export AMD_LOG_LEVEL=4
##DEBUG_FLAGS=-DCMAKE_FIND_DEBUG_MODE=ON
#DEBUG_FLAGS=

CXXFLAGS := $(subst -fstack-protector-strong,-Xarch_host -fstack-protector-strong,$(CXXFLAGS))
CXXFLAGS := $(subst -fcf-protection,-Xarch_host -fcf-protection,$(CXXFLAGS))

# clang-rocm's fork does not bundle its own OpenMP headers (the package is
# deliberately tied to libomp-22, from the distro's generic llvm-22, for
# that); some translation units use -fopenmp=libomp and need omp.h. Stage a
# symlink to omp.h alone in a private directory rather than adding llvm-22's
# whole clang resource dir to the include path, which would leak llvm-22's
# own compiler-internal headers (e.g. __clang_hip_math.h) ahead of
# clang-rocm's matching built-ins and cause missing-type errors from a
# mismatched pair of compiler-internal headers.
OMP_H := /usr/lib/llvm-22/lib/clang/22/include/omp.h
OMP_STAGING_DIR := $(CURDIR)/debian/.omp-include
CXXFLAGS += -isystem $(OMP_STAGING_DIR)
CFLAGS += -isystem $(OMP_STAGING_DIR)

ifdef CCACHE_DIR
CMAKE_CACHE_FLAGS = \
        -D CMAKE_C_COMPILER_LAUNCHER=ccache -D CMAKE_CXX_COMPILER_LAUNCHER=ccache
CCACHE=ccache
endif

# LP_ENV flag: 1 for the Launchpad environment (default), 0 for others.
# TensileCreateLibrary's "Generating assembly kernels" phase hangs when run
# with more than one worker on the constrained Launchpad builders. Force a
# single Tensile worker (TENSILELITE_BUILD_PARALLEL_LEVEL=1) to run it serially
# and complete reliably. Set LP_ENV=0 on a big, well-resourced machine to use
# all available CPUs.
LP_ENV ?= 1
ifeq ($(LP_ENV),1)
TENSILE_PARALLEL_LEVEL := $(shell expr $$(nproc) / 2)
# Ninja job pools (see d/p/0019) can cap concurrent compile/link jobs to
# avoid OOM on memory-constrained builders. Left blank (uncapped, same as
# the LP_ENV=0 branch below): Launchpad's amd64 builders were bumped to
# 32GB/8 threads, so the caps that used to be needed here no longer are.
COMPILE_JOBS :=
LINK_JOBS :=
else
TENSILE_PARALLEL_LEVEL := $(shell nproc)
COMPILE_JOBS :=
LINK_JOBS :=
endif

CMAKE_FLAGS = \
	-DCMAKE_BUILD_TYPE=Release \
	-DTENSILELITE_BUILD_PARALLEL_LEVEL=$(TENSILE_PARALLEL_LEVEL) \
	$(CMAKE_CACHE_FLAGS)

ifeq ($(LP_ENV),1)
CMAKE_FLAGS += \
	-DHIPBLASLT_PARALLEL_COMPILE_JOBS=$(COMPILE_JOBS) \
	-DHIPBLASLT_PARALLEL_LINK_JOBS=$(LINK_JOBS)
endif

# From TheRock build/logs/hipBLASLt_configure.log
CMAKE_FLAGS += \
	-DHIP_PLATFORM=$(shell hipconfig --platform)

CMAKE_FLAGS += \
	-DUSE_SYSTEM_NANOBIND=ON \
	-DDEBIAN_MUT=ON

# Filter out GPU targets that are broken (gfx942) and unsupported (gfx1030)
# and tensilelite missing gfx950
# tensilelite Unsupported GPU target: gfx803
# tensilelite Unsupported GPU target: gfx900
# tensilelite Unsupported GPU target: gfx906
# tensilelite Unsupported GPU target: gfx1010
ifeq ($(LP_ENV),1)
GPU_TARGETS_FILTERED := "gfx1151;gfx1201;gfx908;gfx1150;gfx1200;gfx1101;gfx1100"
else
GPU_TARGETS_FILTERED := "$(shell rocm-target-arch --sep "\n" | grep -vE "(gfx803|gfx9-generic|gfx9-4-generic|gfx900|gfx906|gfx942|gfx950|gfx10-1-generic|gfx1010|gfx10-3-generic|gfx1030|gfx11-generic|gfx12-generic)" | tr '\n' ';' | sed 's/;$$//')"
endif

CMAKE_FLAGS += \
	-DGPU_TARGETS=$(GPU_TARGETS_FILTERED)

# rocRoller not packaged yet
CMAKE_FLAGS += \
	-DHIPBLASLT_ENABLE_ROCROLLER=OFF \
	-DHIPBLASLT_USE_ROCROLLER=OFF

# rocTracer not packaged yet
CMAKE_FLAGS += \
	-DHIPBLASLT_ENABLE_MARKER=OFF

# blis/tests FLAGS
CMAKE_FLAGS += \
	-DHIPBLASLT_ENABLE_BLIS=ON \
	-DTENSILELITE_BUILD_TESTING=ON

# blas/lapack FLAGS (cf install.sh)
CMAKE_FLAGS += \
	-DBLAS_LIBRARIES=/usr/lib/$(DEB_HOST_MULTIARCH)/blas/libblas.a \
	-DLAPACK_LIBRARIES='/usr/lib/$(DEB_HOST_MULTIARCH)/lapack/liblapack.a;/usr/lib/$(DEB_HOST_MULTIARCH)/blas/libblas.a' \
	-DBLA_STATIC=ON

ifeq (,$(filter noinsttest,$(DEB_BUILD_PROFILES)))
CMAKE_FLAGS += \
	-DHIPBLASLT_BUILD_TESTING=ON \
	-DCMAKE_CXX_FLAGS="$(CXXFLAGS) -I/usr/src/googletest/googlemock/include" \
	-DCMAKE_C_FLAGS="$(CFLAGS) -I/usr/src/googletest/googlemock/include"
endif

%:
	dh $@ -Scmake+ninja

override_dh_auto_configure-arch:
	mkdir -p $(OMP_STAGING_DIR)
	ln -sf $(OMP_H) $(OMP_STAGING_DIR)/omp.h
	dh_auto_configure -- $(CMAKE_FLAGS)

# Single-threaded Tensile kernel generation (LP_ENV=1, see above) is slow and
# emits little output for long stretches. On Launchpad, wrap the arch build in
# a heartbeat that prints every 5 minutes so the build is not treated as
# inactive. For LP_ENV=0 (fast multi-core build) run the build unwrapped.
override_dh_auto_build-arch:
ifeq ($(LP_ENV),1)
	@echo "=== Starting build with heartbeat monitoring ==="
	@(while true; do echo "#"; sleep 30m; done) & HB_PID=$$!; \
	trap "kill $$HB_PID 2>/dev/null || true; echo ''; echo '=== Heartbeat stopped ==='" EXIT; \
	set -e; \
	dh_auto_build -a; \
	echo "=== Build completed successfully ==="
else
	dh_auto_build -a
endif

override_dh_strip-arch:
	dh_strip --exclude=.hsaco --exclude=.co --no-automatic-dbgsym

override_dh_dwz-arch:

override_dh_auto_test-arch:
ifeq (,$(filter noinsttest,$(DEB_BUILD_OPTIONS)))
	set -e \
	; if [ -r /dev/kfd ] \
	; then obj-$(DEB_HOST_MULTIARCH)/clients/hipblaslt-test \
	    --gtest_filter="*quick*" \
	; else echo "W: /dev/kfd unreadable: no available AMD GPU." \
	;      echo "W: tests skipped." \
	; fi
endif

# .o files take up a lot of disk space across the multi-arch build, which
# causes Launchpad builders to fail with "No space left on device". They
# are no longer needed by the time dh_auto_install runs (everything needed
# is already linked into libraries/binaries).
execute_before_dh_auto_install-arch:
	find $(CURDIR) -name '*.o' -delete

execute_before_dh_install-arch:

override_dh_auto_configure-indep:

execute_after_dh_auto_build-indep: export http_proxy=127.0.0.1:9
execute_after_dh_auto_build-indep: export https_proxy=127.0.0.1:9
execute_after_dh_auto_build-indep:
ifeq (,$(filter nodoc,$(DEB_BUILD_OPTIONS)))
	rocm-docs-build
	# Fix privacy
	find build/html -name data-type-support.html \
		-exec perl -ni -e 'print unless /cdn\.jsdelivr\.net/' {} \;
	find build/html -name what-is-hipBLASLt.html \
		-exec perl -ni -e 'print unless /cdn\.jsdelivr\.net/' {} \;
	# Remove build path in some html files
	find build/html -type f -name '*.html' -exec sed -i 's|$(CURDIR)/||g' {} +
endif

override_dh_auto_test-indep:

override_dh_auto_install-indep:

execute_before_dh_clean:
	rm -rf $(OMP_STAGING_DIR)

LIB=libhipblaslt1
LIBV=$(shell head -1  debian/changelog | sed 's/.*(\(.*\)-.*).*$$/\1/')
REV=$(shell head -1  debian/changelog | sed 's/.*(.*-\(.*\)).*$$/\1/')
gensymbols:
	rm -rf /tmp/libhipblaslt1
	dpkg-deb -x ../$(LIB)_$(LIBV)-$(REV)_amd64.deb /tmp/$(LIB)
	dpkg-gensymbols -v$(LIBV) -p$(LIB) -P/tmp/$(LIB) -Odebian/$(LIB).symbols
	#Add "* Build-Depends-Package: libhipblaslt-dev" on line 2
	#sed -i '2i * Build-Depends-Package: libhipblaslt-dev' debian/$(LIB).symbols

override_dh_gencontrol:
	dh_gencontrol -- -Vrocm:GPU-Architecture=$(subst ;, ,$(GPU_TARGETS_FILTERED))
