Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 5 additions & 0 deletions projects/hiptensor/CHANGELOG.md
Original file line number Diff line number Diff line change
Expand Up @@ -2,6 +2,11 @@

Full documentation for hipTensor is available at [rocm.docs.amd.com/projects/hiptensor](https://rocm.docs.amd.com/projects/hipTensor/en/latest/index.html).

## hipTensor 2.4.0 for ROCm 7.15

### Added
* Added a host-only stub `libhiptensor`, built automatically when `GPU_TARGETS` resolves to an empty target list (or explicitly with `-DHIPTENSOR_DISABLE_DEVICE=ON`), so the package, headers, and CMake config remain available and every API call returns `HIPTENSOR_STATUS_NOT_SUPPORTED` instead of failing to link.

## hipTensor 2.3.0 for ROCm 7.14

### Added
Expand Down
68 changes: 50 additions & 18 deletions projects/hiptensor/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -74,6 +74,7 @@ set(CMAKE_CXX_FLAGS_MINSIZEREL "${CMAKE_CXX_FLAGS_MINSIZEREL} -O2 -DNDEBUG") # c

# Project options
if(CMAKE_PROJECT_NAME STREQUAL "hiptensor")
option(HIPTENSOR_DISABLE_DEVICE "Build hipTensor without device backend." OFF)
option(HIPTENSOR_BUILD_TESTS "Build hiptensor tests" ON)
option(HIPTENSOR_BUILD_SAMPLES "Build hiptensor samples" ON)
option(HIPTENSOR_BUILD_COMPRESSED_DBG "Enable compressed debug symbols" ON)
Expand Down Expand Up @@ -150,15 +151,30 @@ endif()
# default cached variable GPU_TARGET to "gfx900;gfx906;gfx908", where not all archs are compatible with MFMA instructions
#
# By rule, once cached variable is set, it cannot be overridden unless we use the FORCE option
#
# Three cases are distinguished here:
# * GPU_TARGETS / AMDGPU_TARGETS unset -> fall back to DEFAULT_GPU_TARGETS
# * set to a non-empty list -> honor it
# * explicitly set to an empty string -> no supported architecture is available for this configuration;
# build with stub implementation.
if(GPU_TARGETS)
set(GPU_TARGETS "${GPU_TARGETS}" CACHE STRING "List of specific machine types for library to target")
elseif(AMDGPU_TARGETS)
set(GPU_TARGETS "${AMDGPU_TARGETS}" CACHE STRING "List of specific machine types for library to target")
message(DEPRECATION "AMDGPU_TARGETS use is deprecated. Use GPU_TARGETS.")
elseif(DEFINED GPU_TARGETS OR DEFINED AMDGPU_TARGETS)
# When hipTensor is configured for an architecture set that resolves to no GPU
# targets, the device libraries and their Composable Kernel dependency cannot
# be built. This option is turned on automatically so the library is still built
# with a stub implementation where HIPTENSOR_STATUS_NOT_SUPPORTED is returned
# for any call to device operations.
set(HIPTENSOR_DISABLE_DEVICE ON)
message(WARNING "GPU_TARGETS is empty; hipTensor stub libraries will be built. No device operation supported.")
else()
set(GPU_TARGETS "${DEFAULT_GPU_TARGETS}" CACHE STRING "List of specific machine types for library to target")
endif()
message(STATUS "GPU_TARGETS=${GPU_TARGETS}")
message(STATUS "HIPTENSOR_DISABLE_DEVICE=${HIPTENSOR_DISABLE_DEVICE}")

if(HIPTENSOR_DEFAULT_STRIDES_COL_MAJOR)
add_compile_definitions(HIPTENSOR_DEFAULT_STRIDES_COL_MAJOR=1)
Expand Down Expand Up @@ -204,30 +220,46 @@ endif()
# disable warning about 'amdgpu-waves-per-eu'
add_compile_options(-Wno-pass-failed)

find_package(composable_kernel 1.0.0 REQUIRED COMPONENTS device_contraction_operations device_reduction_operations device_other_operations)
rocm_package_add_dependencies("composable_kernel >= 1.0.0" COMPONENT tests)
# Define the stub helpers
add_subdirectory(library/stub)

# Configure library build
add_subdirectory(library/src)
if(NOT HIPTENSOR_DISABLE_DEVICE)
find_package(composable_kernel 1.0.0 REQUIRED COMPONENTS device_contraction_operations device_reduction_operations device_other_operations)
rocm_package_add_dependencies("composable_kernel >= 1.0.0" COMPONENT tests)

# Configure testing setup
if(HIPTENSOR_BUILD_TESTS OR HIPTENSOR_BUILD_SAMPLES)
enable_testing()
rocm_package_setup_component(clients)
endif()
# Configure library build
add_subdirectory(library/src)

# Configure tests build
if(HIPTENSOR_BUILD_TESTS)
rocm_package_setup_client_component(tests)
add_subdirectory(test)
endif()
# Configure testing setup
if(HIPTENSOR_BUILD_TESTS OR HIPTENSOR_BUILD_SAMPLES)
enable_testing()
rocm_package_setup_component(clients)
endif()

# Configure clients build
if(HIPTENSOR_BUILD_SAMPLES)
rocm_package_setup_client_component(samples)
add_subdirectory(samples)
# Configure tests build
if(HIPTENSOR_BUILD_TESTS)
rocm_package_setup_client_component(tests)
add_subdirectory(test)
endif()

# Configure clients build
if(HIPTENSOR_BUILD_SAMPLES)
rocm_package_setup_client_component(samples)
add_subdirectory(samples)
endif()
else()
# No supported GPU target is available for this configuration, so the device
# library (and its Composable Kernel dependency) cannot be built. Build a
# host-only stub library that implements the full public API but reports
# HIPTENSOR_STATUS_NOT_SUPPORTED for every entry point. This keeps the
# package and its headers available so downstream find_package(hiptensor)
# and linking keep working, while giving callers a well-defined runtime error.
create_hiptensor_stub_library()
endif()

# Guarantee, at build time, that the stub implements the entire public API.
hiptensor_add_stub_api_coverage_check()

# Versioning via rocm-cmake
set (VERSION_STRING "2.3.0")
rocm_setup_version(VERSION ${VERSION_STRING})
Expand Down
1 change: 1 addition & 0 deletions projects/hiptensor/docs/conceptual/programmers-guide.rst
Original file line number Diff line number Diff line change
Expand Up @@ -140,6 +140,7 @@ The ``library`` directory contains the following include and source files:
- ``library/src/reduction/``: Source files for core initialization and management of reduction module.
- ``library/src/reduction/device``: Source files for composable kernel backend reduction instances.
- ``library/src/include``: Infrastructure support for backend and logging management.
- ``library/stub/``: Host-only stub implementation of the public API, built as ``libhiptensor`` when no GPU target is available (see :ref:`installation`). Every entry point returns ``HIPTENSOR_STATUS_NOT_SUPPORTED``. The stub sources are auto-generated from the public header, and a configure-time coverage check keeps them in sync with the API.

``samples`` directory
^^^^^^^^^^^^^^^^^^^^^^^^^^
Expand Down
18 changes: 18 additions & 0 deletions projects/hiptensor/docs/install/installation.rst
Original file line number Diff line number Diff line change
Expand Up @@ -196,6 +196,9 @@ Here are the available options to build the hipTensor library, with or without c
* - ``HIPTENSOR_INLINE_UNARY_OPS``
- Inline all contraction unary ops for best runtime performance (slower compilation)
- ``OFF``
* - ``HIPTENSOR_DISABLE_DEVICE``
- Build a host-only stub library without the device backend (see `Building a stub library`_). This is enabled automatically when ``GPU_TARGETS`` resolves to an empty target list.
- ``OFF``
* - ``CREATE_TEST_APP_LOCAL_DEPLOY``
- Copy ROCm runtime DLLs next to test binaries so they take precedence over System32 (Windows only)
- ``OFF``
Expand Down Expand Up @@ -517,6 +520,21 @@ The following table highlights the relationships between high-level grouped targ
| |``rank6_reduction_test`` |
+-----------------------------------+---------------------------------------------------------------------------------+

Building a stub library
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^

When ``GPU_TARGETS`` resolves to an empty list, the Composable Kernel device backend cannot be
built. In this case hipTensor builds a host-only stub ``libhiptensor`` that implements the full
public API but returns ``HIPTENSOR_STATUS_NOT_SUPPORTED`` for every call, so the package,
headers, and CMake config stay available for downstream consumers.

The stub is selected automatically when ``GPU_TARGETS`` is empty, or explicitly with
``-DHIPTENSOR_DISABLE_DEVICE=ON``:

.. code-block:: bash

CC=/opt/rocm/bin/amdclang CXX=/opt/rocm/bin/amdclang++ cmake -B <build_dir> . -DGPU_TARGETS=""

Building on Windows
-------------------------------------------

Expand Down
215 changes: 215 additions & 0 deletions projects/hiptensor/library/stub/CMakeLists.txt
Original file line number Diff line number Diff line change
@@ -0,0 +1,215 @@
###############################################################################
#
# MIT License
#
# Copyright (C) 2023-2026 Advanced Micro Devices, Inc. All rights reserved.
#
# Permission is hereby granted, free of charge, to any person obtaining a copy
# of this software and associated documentation files (the "Software"), to deal
# in the Software without restriction, including without limitation the rights
# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
# copies of the Software, and to permit persons to whom the Software is
# furnished to do so, subject to the following conditions:
#
# The above copyright notice and this permission notice shall be included in
# all copies or substantial portions of the Software.
#
# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
# THE SOFTWARE.
#
###############################################################################

# Host-only stub of the hipTensor public API.
#
# The stub serves two purposes:
#
# 1. On architectures with no usable GPU target (HIPTENSOR_DISABLE_DEVICE=ON),
# it IS the installed libhiptensor: the device kernels and their Composable
# Kernel dependency cannot be built, but the package, headers, and a
# well-defined runtime error stay available for downstream consumers.
#
# 2. In every build (device or stub), it backs a build-time completeness check
# (hiptensor_add_stub_api_coverage_check below) that guarantees the stub
# implements the entire public API regardless of GPU_TARGETS.
#
# The stub sources are:
# * generated: every public API function returning hiptensorStatus_t, emitted
# from the public header by generate_api_stub.py (--mode stub) so it can
# never drift out of sync with the API.
# * hiptensor_stub_special.cpp: the few informational helpers with real
# behavior (non-status return types), hand-written.

include(GenerateExportHeader)

# These functions are invoked from the parent (root) scope after this directory
# is added, where CMAKE_CURRENT_* would point at the project root. Capture this
# directory now so the helpers reference the stub sources regardless of caller.
set(HIPTENSOR_STUB_DIR "${CMAKE_CURRENT_LIST_DIR}" CACHE INTERNAL "hipTensor stub source dir")
set(HIPTENSOR_STUB_BINARY_DIR "${CMAKE_CURRENT_BINARY_DIR}" CACHE INTERNAL "hipTensor stub binary dir")

# Generate the stub implementation source (status-returning functions) from the
# public header, into ${out_var}. Runs the generator now so the file exists at
# configure time, and registers a build-time custom command so header edits
# refresh it on rebuild.
function(_hiptensor_generate_stub_source out_var)
set(_generated "${HIPTENSOR_STUB_BINARY_DIR}/hiptensor_stub_generated.cpp")
set(${out_var} "${_generated}" PARENT_SCOPE)

# Both the stub library and the coverage check use the generated source;
# register the generation rule only once to avoid a duplicate OUTPUT rule.
get_property(_already GLOBAL PROPERTY HIPTENSOR_STUB_SOURCE_GENERATED)
if(_already)
return()
endif()
set_property(GLOBAL PROPERTY HIPTENSOR_STUB_SOURCE_GENERATED TRUE)

find_package(Python3 COMPONENTS Interpreter REQUIRED)

set(_header "${PROJECT_SOURCE_DIR}/library/include/hiptensor/hiptensor.h")
set(_generator "${HIPTENSOR_STUB_DIR}/generate_api_stub.py")

execute_process(
COMMAND ${Python3_EXECUTABLE} ${_generator}
--mode stub --header ${_header} --output ${_generated}
RESULT_VARIABLE _gen_result
)
if(NOT _gen_result EQUAL 0)
message(FATAL_ERROR "Failed to generate hipTensor stub source")
endif()
add_custom_command(
OUTPUT ${_generated}
COMMAND ${Python3_EXECUTABLE} ${_generator}
--mode stub --header ${_header} --output ${_generated}
DEPENDS ${_generator} ${_header}
COMMENT "Generating hipTensor stub source"
VERBATIM
)
endfunction()

# Build the host-only stub shared library used as libhiptensor when device
# support is disabled.
function(create_hiptensor_stub_library)
include(GenerateExportHeader)
_hiptensor_generate_stub_source(_stub_generated)

add_library(hiptensor SHARED
${_stub_generated}
${HIPTENSOR_STUB_DIR}/hiptensor_stub_special.cpp)
add_library(hiptensor::hiptensor ALIAS hiptensor)

target_compile_options(hiptensor PRIVATE ${CLANG_DRIVER_MODE})
target_link_options(hiptensor PRIVATE ${CLANG_DRIVER_MODE})
target_compile_definitions(hiptensor PRIVATE hiptensor_EXPORTS)
target_include_directories(hiptensor
PUBLIC
$<BUILD_INTERFACE:${PROJECT_SOURCE_DIR}/library/include>
$<BUILD_INTERFACE:${PROJECT_BINARY_DIR}/include>
$<INSTALL_INTERFACE:include>
)
# hiptensor.h includes "internal/hiptensor-export.h" relative to the
# build/include/hiptensor directory, matching the real library build.
target_include_directories(hiptensor PRIVATE ${PROJECT_BINARY_DIR}/include/hiptensor)
target_link_libraries(hiptensor PUBLIC hip::host)
set_target_properties(hiptensor PROPERTIES POSITION_INDEPENDENT_CODE ON)

set(hiptensor_SOVERSION 0.1)
rocm_set_soversion(hiptensor ${hiptensor_SOVERSION})

# hiptensor.h includes an export header normally produced by the device
# build. Generate it here so the installed headers remain self-contained.
generate_export_header(hiptensor
EXPORT_FILE_NAME ${PROJECT_BINARY_DIR}/include/hiptensor/internal/hiptensor-export.h)

rocm_install_targets(
TARGETS hiptensor
INCLUDE
${PROJECT_SOURCE_DIR}/library/include
${PROJECT_BINARY_DIR}/include
)
endfunction()

# Add a build-time check that the stub implements the entire public API.
#
# A coverage translation unit is generated from the public header (one address
# reference per HIPTENSOR_EXPORT function) and linked together with the stub
# sources into a tiny, non-installed host executable. If the stub is missing any
# public API function, that reference is an unresolved symbol and the link --
# hence the build -- fails. Because both the coverage test and the generated stub
# are refreshed from the header, a newly added public API is covered
# automatically. This is independent of GPU_TARGETS and of the device build, so
# it also guards the stub in device builds where the stub is not otherwise
# compiled.
#
# Self-contained: it compiles the stub sources itself and writes its own export
# header, so it has no ordering dependency on the device/stub library.
function(hiptensor_add_stub_api_coverage_check)
if(TARGET hiptensor_stub_api_coverage)
return()
endif()

find_package(Python3 COMPONENTS Interpreter REQUIRED)

set(_header "${PROJECT_SOURCE_DIR}/library/include/hiptensor/hiptensor.h")
set(_generator "${HIPTENSOR_STUB_DIR}/generate_api_stub.py")
set(_generated "${HIPTENSOR_STUB_BINARY_DIR}/hiptensor_stub_api_coverage.cpp")
set(_private_inc "${HIPTENSOR_STUB_BINARY_DIR}/stub_coverage_include")

# hiptensor.h unconditionally includes "internal/hiptensor-export.h", which
# is normally produced by generate_export_header() for a library target.
# This check links an executable, so write a minimal export header directly:
# the symbol-visibility macros only need to be defined for a host compile.
file(WRITE "${_private_inc}/hiptensor/internal/hiptensor-export.h"
"#ifndef HIPTENSOR_EXPORT_H
#define HIPTENSOR_EXPORT_H
#define HIPTENSOR_EXPORT
#define HIPTENSOR_NO_EXPORT
#endif // HIPTENSOR_EXPORT_H
")

# Generate the coverage source now (so it exists at configure) and again at
# build time when the header changes.
execute_process(
COMMAND ${Python3_EXECUTABLE} ${_generator} --header ${_header} --output ${_generated}
RESULT_VARIABLE _gen_result
)
if(NOT _gen_result EQUAL 0)
message(FATAL_ERROR "Failed to generate hipTensor stub API coverage check")
endif()
add_custom_command(
OUTPUT ${_generated}
COMMAND ${Python3_EXECUTABLE} ${_generator} --header ${_header} --output ${_generated}
DEPENDS ${_generator} ${_header}
COMMENT "Generating hipTensor stub API coverage check"
VERBATIM
)

# Link the coverage test against the same stub sources the library uses, so a
# missing implementation fails the link.
_hiptensor_generate_stub_source(_stub_generated)
add_executable(hiptensor_stub_api_coverage
${_generated}
${_stub_generated}
${HIPTENSOR_STUB_DIR}/hiptensor_stub_special.cpp)
# Exclude from the default build set; depended on explicitly below so it
# still builds with the library but does not get installed or packaged.
set_target_properties(hiptensor_stub_api_coverage PROPERTIES EXCLUDE_FROM_ALL TRUE)
target_compile_options(hiptensor_stub_api_coverage PRIVATE ${CLANG_DRIVER_MODE})
target_link_options(hiptensor_stub_api_coverage PRIVATE ${CLANG_DRIVER_MODE})
target_compile_definitions(hiptensor_stub_api_coverage PRIVATE hiptensor_EXPORTS)
target_include_directories(hiptensor_stub_api_coverage PRIVATE
${PROJECT_SOURCE_DIR}/library/include
${_private_inc}
${_private_inc}/hiptensor)
target_link_libraries(hiptensor_stub_api_coverage PRIVATE hip::host)

# Tie the check into the default build so it runs on every build of the
# library, turning a missing stub implementation into a build failure.
if(TARGET hiptensor)
add_dependencies(hiptensor hiptensor_stub_api_coverage)
endif()
endfunction()
Loading
Loading