From e0b0a896381142d175baf03f586392a50d17273e Mon Sep 17 00:00:00 2001 From: mousdahl-amd Date: Tue, 23 Jun 2026 15:23:37 -0700 Subject: [PATCH 1/2] Final bit of groundwork before implementing graph wrapper --- .../hipdnn_compatibility/cudnn/cudnn.h | 14 +- .../cudnn/cudnn_frontend.h | 3 + .../cudnn/cudnn_frontend/graph_helpers.h | 119 +++++++++++++++++ .../cudnn/cudnn_frontend/graph_properties.h | 52 ++++++++ .../cudnn/cudnn_frontend_utils.h | 65 +++++++++ projects/hipdnn/frontend/tests/CMakeLists.txt | 3 + .../frontend/tests/TestCudnnShimError.cpp | 96 ++++++++++++++ .../frontend/tests/TestCudnnShimTensor.cpp | 124 ++++++++++++++++++ .../frontend/tests/TestCudnnShimV9Tu.cpp | 59 +++++++++ 9 files changed, 529 insertions(+), 6 deletions(-) create mode 100644 projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend/graph_helpers.h create mode 100644 projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend/graph_properties.h create mode 100644 projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend_utils.h create mode 100644 projects/hipdnn/frontend/tests/TestCudnnShimError.cpp create mode 100644 projects/hipdnn/frontend/tests/TestCudnnShimTensor.cpp create mode 100644 projects/hipdnn/frontend/tests/TestCudnnShimV9Tu.cpp diff --git a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn.h b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn.h index a8997f47eac7..9810ee3beddd 100644 --- a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn.h +++ b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn.h @@ -52,14 +52,16 @@ static_assert(std::is_same_v, // status enum is available without the C entry points — this lets // `detail/status_translation.h` be self-contained. -// NOTE: This stub intentionally declares only the C-API types the *implemented* -// entry points use — `cudnnHandle_t` and `cudnnStatus_t`. The remaining v9 -// C-API enums named in RFC 0012 §4.7 (`cudnnDataType_t`, `cudnnTensorFormat_t`, +// NOTE: This stub intentionally declares only the C-API types the v9 graph API +// actually references — `cudnnHandle_t` and `cudnnStatus_t`. The other C-API +// enums named in RFC 0012 §4.7 (`cudnnDataType_t`, `cudnnTensorFormat_t`, // `cudnnConvolutionMode_t`, `cudnnReduceTensorOp_t`, `cudnnNormFwdPhase_t`, // `cudnnBackendHeurMode_t`, `cudnnBackendNumericalNote_t`, -// `cudnnBackendBehaviorNote_t`, `cudnnBackendDescriptorType_t`) land with the -// type-mapping work, where their values/aliasing are verified against upstream -// rather than stubbed here. +// `cudnnBackendBehaviorNote_t`, `cudnnBackendDescriptorType_t`) are deliberately +// NOT declared: no public v9 graph signature names them. The v9 surface uses the +// FE-namespace enums (`cudnn_frontend::DataType_t`, …) — which the shim *aliases* +// to the hipDNN types (see `cudnn_frontend_utils.h`), not these C-API enums. +// They will be added only if a future consumer surface actually references one. // Status translation between the cuDNN and hipDNN enum families. Included after // the C-API types above (it needs cudnnStatus_t) and before the entry points diff --git a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend.h b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend.h index ce989ca220c3..da897935f0bd 100644 --- a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend.h +++ b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend.h @@ -34,4 +34,7 @@ #pragma once #include +#include +#include +#include #include diff --git a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend/graph_helpers.h b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend/graph_helpers.h new file mode 100644 index 000000000000..00325e58691b --- /dev/null +++ b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend/graph_helpers.h @@ -0,0 +1,119 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT +// +// Portions derived from NVIDIA cuDNN frontend (include/cudnn_frontend/graph_helpers.h +// and include/cudnn_frontend_Logging.h), used under the MIT license. The error +// type names, error-code enum, and error/log macro names below intentionally +// mirror cuDNN frontend so hipified v9 consumers compile unchanged. + +/** + * @file graph_helpers.h + * @brief Error-type aliases and error/log macros for the hipDNN cuDNN shim. + * + * cuDNN frontend's `error_object` / `error_t` / `error_code_t` (defined upstream + * in `cudnn_frontend/graph_helpers.h`) are structurally identical to + * hipDNN's `Error` / `ErrorCode`, and hipDNN's error-code enum is a + * name-superset of cuDNN's (RFC 0012 §4.6). This header therefore **aliases** + * the hipDNN types into `` with `using` declarations rather than + * wrapping them — `cudnn_frontend::error_t` *is* `hipdnn_frontend::error_t`, so + * an error returned by a wrapped hipDNN call flows out unchanged, with no cast + * between enum families (RFC 0012 §4.3). + * + * It also re-exports the cuDNN frontend error/log macros + * (`CHECK_CUDNN_FRONTEND_ERROR`, `RETURN_CUDNN_FRONTEND_ERROR_IF`, + * `CUDNN_FE_LOG*`) — PyTorch's `AT_CUDNN_FRONTEND_CHECK` expands to them. The + * log macros delegate to hipDNN's frontend logging (off by default, gated by + * `HIPDNN_LOG_LEVEL`); the full cuDNN-FE log env-var bridge (RFC 0012 §4.6) + * lands in a later phase. + * + * @note Internal-to-shim; pulled in by the umbrella `cudnn_frontend.h`. + */ + +#pragma once + +#include +#include + +namespace hipdnn_frontend::compatibility::cudnn_frontend +{ + +// Error types aliased 1:1 (RFC 0012 §4.6). hipDNN's `error_code_t` is a +// name-superset of cuDNN's 16 codes, and `error_object` / `error_t` share +// cuDNN FE's public surface (code, err_msg, get_code/get_message/is_good/ +// is_bad/operator==/!=), so the alias is exact. +using hipdnn_frontend::error_code_t; +using hipdnn_frontend::error_object; +using hipdnn_frontend::error_t; + +} // namespace hipdnn_frontend::compatibility::cudnn_frontend + +// =========================================================================== +// Error/log macros (RFC 0012 §4.6) — re-exported with the cuDNN FE names. +// =========================================================================== +// These are global (un-namespaced), exactly as upstream ships them, because +// consumer code (e.g. PyTorch's AT_CUDNN_FRONTEND_CHECK) uses the bare names. +// `error_t` / `error_code_t` referenced inside expand against whatever +// `cudnn_frontend` alias the consumer has in scope, matching upstream behavior. + +// The log macros forward their `<<`-streamed argument straight into hipDNN's +// frontend INFO logger. HIPDNN_FE_LOG_INFO gates on HIPDNN_LOG_LEVEL (off by +// default) and accepts a stream expression, so chained `<<` arguments work. +#ifndef CUDNN_FE_LOG +#define CUDNN_FE_LOG(X) HIPDNN_FE_LOG_INFO(X) +#endif + +// The X argument is a stream expression: it must stay un-parenthesized so it +// chains off the leftmost LogStream operand (parenthesizing it would force a +// `const char* << ...` evaluation and break compilation). Same NOLINT pattern +// the SDK uses for its own streaming log macros. +#ifndef CUDNN_FE_LOG_LABEL +#define CUDNN_FE_LOG_LABEL(X) \ + HIPDNN_FE_LOG_INFO("[cudnn_frontend] " << X) // NOLINT(bugprone-macro-parentheses) +#endif + +#ifndef CUDNN_FE_LOG_LABEL_ENDL +#define CUDNN_FE_LOG_LABEL_ENDL(X) \ + HIPDNN_FE_LOG_INFO("[cudnn_frontend] " << X) // NOLINT(bugprone-macro-parentheses) +#endif + +#ifndef CUDNN_FE_LOG_BANNER +#define CUDNN_FE_LOG_BANNER(X) \ + HIPDNN_FE_LOG_INFO("[cudnn_frontend] === " << X << " ===") // NOLINT(bugprone-macro-parentheses) +#endif + +/// @brief Evaluate an expression returning `error_t`; on `is_bad()`, log and +/// propagate it. Mirrors cuDNN FE's `CHECK_CUDNN_FRONTEND_ERROR`. +#ifndef CHECK_CUDNN_FRONTEND_ERROR +#define CHECK_CUDNN_FRONTEND_ERROR(x) \ + do \ + { \ + if(auto retval = (x); retval.is_bad()) \ + { \ + CUDNN_FE_LOG_LABEL_ENDL("ERROR: " << #x << " at " << __FILE__ << ":" << __LINE__); \ + return retval; \ + } \ + } while(0) +#endif + +/// @brief If `cond`, log and `return {retval, message}`. Mirrors cuDNN FE's +/// `RETURN_CUDNN_FRONTEND_ERROR_IF`. +#ifndef RETURN_CUDNN_FRONTEND_ERROR_IF +#define RETURN_CUDNN_FRONTEND_ERROR_IF(cond, retval, message) \ + do \ + { \ + if(cond) \ + { \ + if((retval) == error_code_t::OK) \ + { \ + CUDNN_FE_LOG_LABEL("INFO: "); \ + } \ + else \ + { \ + CUDNN_FE_LOG_LABEL("ERROR: "); \ + } \ + CUDNN_FE_LOG((message) << ". because (" << #cond ") at " << __FILE__ << ":" \ + << __LINE__ << "\n"); \ + return {retval, message}; \ + } \ + } while(0) +#endif diff --git a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend/graph_properties.h b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend/graph_properties.h new file mode 100644 index 000000000000..4c68e49fd0ef --- /dev/null +++ b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend/graph_properties.h @@ -0,0 +1,52 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT +// +// Portions derived from NVIDIA cuDNN frontend +// (include/cudnn_frontend/graph_properties.h), used under the MIT license. The +// type name below intentionally mirrors cuDNN frontend so hipified v9 consumers +// compile unchanged. + +/** + * @file graph_properties.h + * @brief Graph attribute-type aliases for the hipDNN cuDNN-compatibility shim. + * + * cuDNN frontend declares its v9 graph attribute structs (`Tensor_attributes`, + * the per-node `*_attributes`, …) in `cudnn_frontend/graph_properties.h`. This + * header mirrors that filename and brings the matching hipDNN types into + * `::graph` by **aliasing** them with `using` declarations — zero + * overhead, no re-declaration, and no shim-side state (RFC 0012 §4.4.1). + * + * `cudnn_frontend::graph::Tensor_attributes` therefore *is* + * `hipdnn_frontend::graph::TensorAttributes`: a tensor configured through the + * shim flows into a wrapped hipDNN graph with no conversion, and tensor UID + * handling stays entirely on the hipDNN side (`assignUnsetTensorUids()` at + * build). The shim adds no per-tensor identity map or UID allocator + * (RFC 0012 §4.4.1, §5.3). + * + * @par Setter coverage (RFC 0012 §4.4.1) + * hipDNN's `TensorAttributes` publishes the cuDNN-named chained setters the v9 + * graph API uses: `set_dim`, `set_stride`, `set_data_type`, `set_uid`, + * `set_is_virtual`, `set_output`, `set_name` — aliased here 1:1. Pass-by-value + * tensors are expressed through hipDNN's typed `set_value(scalar)` / + * scalar-constructor surface rather than a separate `set_is_pass_by_value(bool)` + * flag. cuDNN FE's `set_is_pass_by_value` and `set_ragged_offset` have no + * hipDNN counterpart yet (no ragged-tensor support); they are intentionally + * out of this alias's scope and are added on the hipDNN side, then surfaced + * here, when their consuming nodes land (RFC 0012 §7.x). + * + * @note Internal-to-shim; pulled in by the umbrella `cudnn_frontend.h`. + */ + +#pragma once + +#include + +namespace hipdnn_frontend::compatibility::cudnn_frontend::graph +{ + +// Tensor attributes aliased 1:1 (RFC 0012 §4.4.1). hipDNN publishes both the +// class name and cuDNN's `Tensor_attributes` spelling (a typedef); both are +// aliased so consumer code using either name resolves through the shim. +using hipdnn_frontend::graph::Tensor_attributes; + +} // namespace hipdnn_frontend::compatibility::cudnn_frontend::graph diff --git a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend_utils.h b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend_utils.h new file mode 100644 index 000000000000..db4d05968c29 --- /dev/null +++ b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend_utils.h @@ -0,0 +1,65 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT +// +// Portions derived from NVIDIA cuDNN frontend (include/cudnn_frontend_utils.h), +// used under the MIT license. The enum names below intentionally mirror the +// cuDNN frontend FE-namespace enums so hipified v9 consumers compile unchanged. + +/** + * @file cudnn_frontend_utils.h + * @brief FE-namespace enum aliases for the hipDNN cuDNN-compatibility shim. + * + * cuDNN frontend uses two parallel enum families (RFC 0012 §4.3): the C-API + * enums from `` and the FE-namespace C++ enums in `namespace + * cudnn_frontend` (`DataType_t`, `PointwiseMode_t`, …). The v9 graph API + * signatures use the **FE-namespace** family; this header makes those names + * available under `` by **aliasing** the matching hipDNN types with + * `using` declarations — zero overhead, no re-declaration, and no numeric cast + * between enum families (RFC 0012 §4.3, §4.3.1). + * + * hipDNN already publishes these enums as cuDNN-named `_t` typedefs (see + * `hipdnn_frontend/Types.hpp`), each a name-superset of its cuDNN counterpart, + * so the alias holds. + * + * @note Internal-to-shim; pulled in by the umbrella `cudnn_frontend.h`. + */ + +#pragma once + +#include + +namespace hipdnn_frontend::compatibility::cudnn_frontend +{ + +// FE-namespace enums that hipDNN publishes 1:1 (RFC 0012 §4.3). Aliased, not +// re-declared, so `cudnn_frontend::DataType_t` *is* `hipdnn_frontend::DataType_t`. +using hipdnn_frontend::AttentionImplementation_t; +using hipdnn_frontend::BehaviorNote_t; +using hipdnn_frontend::BuildPlanPolicy_t; +using hipdnn_frontend::ConvolutionMode_t; +using hipdnn_frontend::DataType_t; +using hipdnn_frontend::DiagonalAlignment_t; +using hipdnn_frontend::HeurMode_t; +using hipdnn_frontend::NormFwdPhase_t; +using hipdnn_frontend::PaddingMode_t; +using hipdnn_frontend::PointwiseMode_t; +using hipdnn_frontend::ReductionMode_t; +using hipdnn_frontend::ResampleMode_t; + +// Other cuDNN FE-namespace enums (NumericalNote_t, NormMode_t, RngDistribution_t, +// DescriptorType_t, MoeGroupedMatmulMode_t, TensorReordering_t, ReshapeMode_t) +// are intentionally NOT aliased here: hipDNN does not yet publish them, and the +// nodes that use them are out of scope until later phases (RFC 0012 §7.3, §7.4). +// They are added on the hipDNN side and aliased when their node lands. + +} // namespace hipdnn_frontend::compatibility::cudnn_frontend + +// The `graph` sub-namespace is populated by the `cudnn_frontend/*` headers the +// umbrella pulls in: `graph_properties.h` aliases the attribute types +// (`Tensor_attributes`, …). The `Graph` composition wrapper and the remaining +// per-node `*_attributes` aliases land in later phases (RFC 0012 §7.2). +// Declaring the namespace here keeps this header self-contained even when +// included on its own. +namespace hipdnn_frontend::compatibility::cudnn_frontend::graph +{ +} // namespace hipdnn_frontend::compatibility::cudnn_frontend::graph diff --git a/projects/hipdnn/frontend/tests/CMakeLists.txt b/projects/hipdnn/frontend/tests/CMakeLists.txt index bfbd3e47349e..b6b177c42f88 100644 --- a/projects/hipdnn/frontend/tests/CMakeLists.txt +++ b/projects/hipdnn/frontend/tests/CMakeLists.txt @@ -40,6 +40,9 @@ add_executable( $<$:TestCudnnShimUmbrellaCompiles.cpp> $<$:TestCudnnShimHandle.cpp> $<$:TestCudnnShimStatusTranslationStandalone.cpp> + $<$:TestCudnnShimError.cpp> + $<$:TestCudnnShimTensor.cpp> + $<$:TestCudnnShimV9Tu.cpp> TestDescriptorHelpers.cpp TestDescriptorUnpackHelpers.cpp TestConvolutionWgradNode.cpp diff --git a/projects/hipdnn/frontend/tests/TestCudnnShimError.cpp b/projects/hipdnn/frontend/tests/TestCudnnShimError.cpp new file mode 100644 index 000000000000..24d5968aaf6d --- /dev/null +++ b/projects/hipdnn/frontend/tests/TestCudnnShimError.cpp @@ -0,0 +1,96 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +// error_object construction, code mapping, and message +// preservation through the aliased hipdnn_frontend::Error, plus the re-exported +// error macros. The shim aliases the hipDNN error types 1:1 (no wrapper), so +// these assertions also pin that the alias is exact. Gated behind +// HIPDNN_ENABLE_CUDNN_COMPATIBILITY in the frontend tests CMakeLists. +#include + +#include + +#include + +namespace cfe = hipdnn_frontend::compatibility::cudnn_frontend; + +// The error types must be the *same* types as hipDNN's, not parallel +// re-declarations (RFC 0012 §4.6 — alias, do not wrap). +static_assert(std::is_same_v, + "cudnn_frontend::error_t must alias hipdnn_frontend::error_t"); +static_assert(std::is_same_v, + "cudnn_frontend::error_object must alias hipdnn_frontend::error_object"); +static_assert(std::is_same_v, + "cudnn_frontend::error_code_t must alias hipdnn_frontend::error_code_t"); + +namespace +{ +// Mirror a hipified consumer: bring the shim names into scope so the bare +// macro identifiers (error_code_t, error_t) resolve as they do upstream. +using namespace hipdnn_frontend::compatibility::cudnn_frontend; + +TEST(TestCudnnShimError, DefaultObjectIsOkAndGood) +{ + const error_t e; + EXPECT_TRUE(e.is_good()); + EXPECT_FALSE(e.is_bad()); + EXPECT_EQ(e.get_code(), error_code_t::OK); + EXPECT_TRUE(e == error_code_t::OK); +} + +TEST(TestCudnnShimError, PreservesCodeAndMessage) +{ + const error_object e{error_code_t::GRAPH_NOT_SUPPORTED, "no engine accepted graph"}; + EXPECT_TRUE(e.is_bad()); + EXPECT_FALSE(e.is_good()); + EXPECT_EQ(e.get_code(), error_code_t::GRAPH_NOT_SUPPORTED); + EXPECT_EQ(e.get_message(), "no engine accepted graph"); + EXPECT_TRUE(e == error_code_t::GRAPH_NOT_SUPPORTED); + EXPECT_TRUE(e != error_code_t::OK); +} + +// Because the shim error type *is* hipDNN's Error, a value produced on one side +// is read on the other with no conversion. +TEST(TestCudnnShimError, AliasInteroperatesWithHipdnnError) +{ + const hipdnn_frontend::error_t native{error_code_t::INVALID_VALUE, "bad value"}; + // Same type on both sides — a cudnn_frontend::error_t reference binds to a + // hipdnn_frontend::error_t with no conversion. + const error_t& aliased = native; + EXPECT_EQ(aliased.get_code(), error_code_t::INVALID_VALUE); + EXPECT_EQ(aliased.get_message(), "bad value"); +} + +// Helpers exercising the re-exported error macros. +error_t failsIfNegative(int x) +{ + RETURN_CUDNN_FRONTEND_ERROR_IF(x < 0, error_code_t::INVALID_VALUE, "x must be >= 0"); + return {}; +} + +error_t checkThenSucceed(int x) +{ + CHECK_CUDNN_FRONTEND_ERROR(failsIfNegative(x)); + return {}; +} + +TEST(TestCudnnShimError, ReturnErrorIfMacro) +{ + auto bad = failsIfNegative(-1); + EXPECT_TRUE(bad.is_bad()); + EXPECT_EQ(bad.get_code(), error_code_t::INVALID_VALUE); + EXPECT_EQ(bad.get_message(), "x must be >= 0"); + + EXPECT_TRUE(failsIfNegative(5).is_good()); +} + +TEST(TestCudnnShimError, CheckErrorMacroPropagates) +{ + auto bad = checkThenSucceed(-1); + EXPECT_TRUE(bad.is_bad()); + EXPECT_EQ(bad.get_code(), error_code_t::INVALID_VALUE); + + EXPECT_TRUE(checkThenSucceed(2).is_good()); +} + +} // namespace diff --git a/projects/hipdnn/frontend/tests/TestCudnnShimTensor.cpp b/projects/hipdnn/frontend/tests/TestCudnnShimTensor.cpp new file mode 100644 index 000000000000..808c278fd86a --- /dev/null +++ b/projects/hipdnn/frontend/tests/TestCudnnShimTensor.cpp @@ -0,0 +1,124 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +// AC for the shim Tensor_attributes (RFC 0012 §4.4.1): a cuDNN-shaped +// `graph::Tensor_attributes` that hipified v9 graph-construction code declares +// and configures against the shim. The shim aliases hipDNN's native +// `graph::TensorAttributes` 1:1 (no wrapper, no shim-side UID allocator or +// identity map), so these assertions both exercise the setter surface and pin +// that the alias is exact. Gated behind HIPDNN_ENABLE_CUDNN_COMPATIBILITY in +// the frontend tests CMakeLists. +#include + +#include + +#include +#include + +namespace cfe = hipdnn_frontend::compatibility::cudnn_frontend; + +// The shim attribute type must be the *same* type as hipDNN's, not a parallel +// re-declaration (RFC 0012 §4.4.1 — alias, do not wrap). Both the class name +// and cuDNN's `Tensor_attributes` spelling resolve to the one hipDNN type. +static_assert( + std::is_same_v, + "cudnn_frontend::graph::Tensor_attributes must alias the hipDNN type"); +static_assert(std::is_same_v, + "the Tensor_attributes and TensorAttributes spellings must be the same type"); + +namespace +{ +// Mirror a hipified consumer (RFC 0012 §4.2 workflow 1): resolve graph symbols +// through the shim namespace. +namespace fe = hipdnn_frontend::compatibility::cudnn_frontend; + +TEST(TestCudnnShimTensor, DefaultConstructs) +{ + const fe::graph::Tensor_attributes tensor; + EXPECT_EQ(tensor.get_uid(), 0); + EXPECT_FALSE(tensor.has_uid()); + EXPECT_EQ(tensor.get_name(), ""); + EXPECT_EQ(tensor.get_data_type(), fe::DataType_t::NOT_SET); + EXPECT_TRUE(tensor.get_dim().empty()); + EXPECT_TRUE(tensor.get_stride().empty()); + EXPECT_FALSE(tensor.get_is_virtual()); +} + +// Every setter named in the acceptance criteria that hipDNN publishes, exercised +// through the shim alias. The setters chain (return *this), exactly as cuDNN FE. +TEST(TestCudnnShimTensor, AllSettersChainAndApply) +{ + fe::graph::Tensor_attributes tensor; + tensor.set_dim({1, 64, 28, 28}) + .set_stride({50176, 784, 28, 1}) + .set_data_type(fe::DataType_t::HALF) + .set_uid(7) + .set_is_virtual(false) + .set_name("input_x"); + + EXPECT_EQ(tensor.get_dim(), std::vector({1, 64, 28, 28})); + EXPECT_EQ(tensor.get_stride(), std::vector({50176, 784, 28, 1})); + EXPECT_EQ(tensor.get_data_type(), fe::DataType_t::HALF); + EXPECT_EQ(tensor.get_uid(), 7); + EXPECT_TRUE(tensor.has_uid()); + EXPECT_FALSE(tensor.get_is_virtual()); + EXPECT_EQ(tensor.get_name(), "input_x"); +} + +// set_output is the cuDNN convenience inverse of set_is_virtual. +TEST(TestCudnnShimTensor, SetOutputIsInverseOfVirtual) +{ + fe::graph::Tensor_attributes tensor; + tensor.set_output(true); + EXPECT_FALSE(tensor.get_is_virtual()); + tensor.set_output(false); + EXPECT_TRUE(tensor.get_is_virtual()); +} + +// UID assignment / clearing forwards to hipDNN's native UID handling; the shim +// keeps no identity map of its own (RFC 0012 §4.4.1, §5.3). +TEST(TestCudnnShimTensor, UidAssignmentForwardsToHipdnn) +{ + fe::graph::Tensor_attributes tensor; + EXPECT_FALSE(tensor.has_uid()); + + tensor.set_uid(42); + EXPECT_EQ(tensor.get_uid(), 42); + EXPECT_TRUE(tensor.has_uid()); + + tensor.clear_uid(); + EXPECT_EQ(tensor.get_uid(), 0); + EXPECT_FALSE(tensor.has_uid()); +} + +// Copying carries the full configuration, including UID identity — there is no +// shim-side per-tensor state that a copy could drop or alias. +TEST(TestCudnnShimTensor, IdentityPreservedAcrossCopies) +{ + fe::graph::Tensor_attributes original; + original.set_dim({2, 3}) + .set_stride({3, 1}) + .set_data_type(fe::DataType_t::FLOAT) + .set_uid(99) + .set_name("t"); + + const fe::graph::Tensor_attributes copyConstructed(original); + EXPECT_EQ(copyConstructed.get_uid(), 99); + EXPECT_TRUE(copyConstructed.has_uid()); + EXPECT_EQ(copyConstructed.get_name(), "t"); + EXPECT_EQ(copyConstructed.get_dim(), std::vector({2, 3})); + EXPECT_EQ(copyConstructed.get_data_type(), fe::DataType_t::FLOAT); + + fe::graph::Tensor_attributes copyAssigned; + copyAssigned = original; + EXPECT_EQ(copyAssigned.get_uid(), 99); + EXPECT_TRUE(copyAssigned.has_uid()); + EXPECT_EQ(copyAssigned.get_name(), "t"); + + // A reference to the shim type binds a hipDNN value with no conversion, + // because they are the same type. + const hipdnn_frontend::graph::TensorAttributes& asNative = original; + EXPECT_EQ(asNative.get_uid(), 99); +} + +} // namespace diff --git a/projects/hipdnn/frontend/tests/TestCudnnShimV9Tu.cpp b/projects/hipdnn/frontend/tests/TestCudnnShimV9Tu.cpp new file mode 100644 index 000000000000..9d46a90d94ce --- /dev/null +++ b/projects/hipdnn/frontend/tests/TestCudnnShimV9Tu.cpp @@ -0,0 +1,59 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +// (RFC 0012 §7.1 Phase-1 exit criterion): a hand-written +// v9-only translation unit that includes ONLY the cuDNN-compatibility shim +// umbrella header compiles and links, using the FE-namespace enums, the error +// types, and the §4.7 C-API types it needs. This is the source-compatibility +// proof for the enum/error/macro surface delivered by this story; the graph +// and attribute wrappers arrive in Phase 2. Gated behind +// HIPDNN_ENABLE_CUDNN_COMPATIBILITY in the frontend tests CMakeLists. +#include + +#include + +// Mirror the consumer hipification step (RFC 0012 §4.2 workflow 1): alias the +// namespace so in-code `cudnn_frontend::` symbols resolve to the shim. +namespace cudnn_frontend = hipdnn_frontend::compatibility::cudnn_frontend; + +// FE-namespace enums resolve through the shim aliases and are the *same* enum +// as the hipDNN type (aliased, not re-declared). +static_assert(std::is_same_v, + "cudnn_frontend::DataType_t must alias the hipDNN enum"); + +namespace +{ +constexpr cudnn_frontend::DataType_t K_IO_TYPE = cudnn_frontend::DataType_t::HALF; +constexpr cudnn_frontend::DataType_t K_COMPUTE_TYPE = cudnn_frontend::DataType_t::FLOAT; +constexpr cudnn_frontend::HeurMode_t K_HEUR_MODE = cudnn_frontend::HeurMode_t::A; + +// error_t / error_code_t resolve through the shim aliases. +cudnn_frontend::error_t makeOk() +{ + return {}; +} + +// §4.7 C-API types come from the stub pulled in by the umbrella. +cudnnStatus_t statusFor(cudnnHandle_t /*handle*/) +{ + return CUDNN_STATUS_SUCCESS; +} + +TEST(TestCudnnShimV9Tu, V9OnlyTranslationUnitCompilesAndLinks) +{ + // The graph sub-namespace resolves even though it is empty at this phase + // (this using-directive would fail to compile if the namespace did not + // exist; Phase 2 will populate it with the Graph/attribute wrappers). + using namespace cudnn_frontend::graph; + + EXPECT_TRUE(makeOk().is_good()); + EXPECT_EQ(K_IO_TYPE, cudnn_frontend::DataType_t::HALF); + EXPECT_EQ(K_COMPUTE_TYPE, cudnn_frontend::DataType_t::FLOAT); + EXPECT_EQ(K_HEUR_MODE, cudnn_frontend::HeurMode_t::A); + EXPECT_EQ(statusFor(nullptr), CUDNN_STATUS_SUCCESS); + EXPECT_EQ(CUDNN_FRONTEND_VERSION, 12400); + + SUCCEED() << "v9-only TU compiled and linked against the shim umbrella."; +} + +} // namespace From 43b1356549d00ecbd989e5daa36468220fb9c186 Mon Sep 17 00:00:00 2001 From: mousdahl-amd Date: Wed, 24 Jun 2026 11:48:08 -0700 Subject: [PATCH 2/2] Trimmed comments to be a little less verbose --- .../hipdnn_compatibility/cudnn/cudnn.h | 71 +++++++------------ .../cudnn/cudnn_frontend.h | 23 ++---- .../cudnn/cudnn_frontend/graph_helpers.h | 33 +-------- .../cudnn/cudnn_frontend/graph_properties.h | 37 +++------- .../cudnn/cudnn_frontend_utils.h | 33 +++------ .../cudnn/cudnn_frontend_version.h | 21 ++---- .../cudnn/cudnn_runtime_version.h | 27 +++---- .../hipdnn_compatibility/cudnn/cudnn_status.h | 8 +-- .../cudnn/detail/status_translation.h | 18 ++--- 9 files changed, 78 insertions(+), 193 deletions(-) diff --git a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn.h b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn.h index 9810ee3beddd..8164de4b84a2 100644 --- a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn.h +++ b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn.h @@ -7,18 +7,14 @@ * @file cudnn.h * @brief Stub C-API header for the hipDNN cuDNN-compatibility shim (v9-only). * - * This is a hand-curated, **v9-only** replacement for NVIDIA's `` - * (RFC 0012 §4.7). It declares only the small set of C-API types that the - * cuDNN frontend v9 graph API refers to in its own method signatures, plus the - * handful of C entry points needed for handle init, stream binding, error - * handling, and version checks (and to build the mirrored samples, §8.3). - * - * Everything here forwards to an existing hipDNN equivalent; the full cuDNN C - * library (convolution-descriptor APIs, etc.) is intentionally out of scope. + * Hand-curated, v9-only replacement for NVIDIA's ``: declares the few + * C-API types the cuDNN frontend v9 graph API names in its signatures, plus the + * C entry points for handle init, stream binding, error handling, and version + * checks. Everything forwards to an existing hipDNN equivalent. * * @note Forwarding goes through `hipdnn_frontend::detail::hipdnnBackend()` — the - * same mockable indirection `hipdnn_frontend/Handle.hpp` uses — so these - * entry points are unit-testable with the in-tree mock backend. + * same mockable indirection `Handle.hpp` uses — so these entry points are + * unit-testable with the in-tree mock backend. */ #pragma once @@ -35,41 +31,29 @@ #include // =========================================================================== -// C-API types (RFC 0012 §4.7) +// C-API types // =========================================================================== /// @brief cuDNN handle, aliased directly to the hipDNN handle type. /// -/// `hipdnnHandle_t` is the global C typedef from `` -/// (`typedef struct hipdnnHandle* hipdnnHandle_t;`), not a member of namespace -/// `hipdnn_frontend`. +/// `hipdnnHandle_t` is the global C typedef from ``, not a +/// member of namespace `hipdnn_frontend`. using cudnnHandle_t = ::hipdnnHandle_t; static_assert(std::is_same_v, - "cudnnHandle_t must alias the hipDNN handle type (RFC 0012 §4.7)"); - -// `cudnnStatus_t` is declared in `cudnn_status.h` (included above) so the -// status enum is available without the C entry points — this lets -// `detail/status_translation.h` be self-contained. - -// NOTE: This stub intentionally declares only the C-API types the v9 graph API -// actually references — `cudnnHandle_t` and `cudnnStatus_t`. The other C-API -// enums named in RFC 0012 §4.7 (`cudnnDataType_t`, `cudnnTensorFormat_t`, -// `cudnnConvolutionMode_t`, `cudnnReduceTensorOp_t`, `cudnnNormFwdPhase_t`, -// `cudnnBackendHeurMode_t`, `cudnnBackendNumericalNote_t`, -// `cudnnBackendBehaviorNote_t`, `cudnnBackendDescriptorType_t`) are deliberately -// NOT declared: no public v9 graph signature names them. The v9 surface uses the -// FE-namespace enums (`cudnn_frontend::DataType_t`, …) — which the shim *aliases* -// to the hipDNN types (see `cudnn_frontend_utils.h`), not these C-API enums. -// They will be added only if a future consumer surface actually references one. - -// Status translation between the cuDNN and hipDNN enum families. Included after -// the C-API types above (it needs cudnnStatus_t) and before the entry points -// below (which use it). Lives under detail/ and is shim-internal. + "cudnnHandle_t must alias the hipDNN handle type"); + +// Only the C-API types the v9 graph API actually references are declared here +// (cudnnHandle_t, plus cudnnStatus_t from cudnn_status.h). Other cuDNN C-API +// enums are intentionally omitted: the v9 graph surface uses the FE-namespace +// enums (DataType_t, …) aliased in cudnn_frontend_utils.h, not the C-API ones. + +// Status translation. Included after cudnnStatus_t above and before the entry +// points below, which use it. #include // =========================================================================== -// C entry points (RFC 0012 §4.7) — forward to the hipDNN backend +// C entry points — forward to the hipDNN backend // =========================================================================== extern "C" { @@ -120,10 +104,8 @@ inline size_t cudnnGetVersion(void) // =========================================================================== // create_cudnn_handle() convenience helper // =========================================================================== -// Mirrors the helper NVIDIA ships in its sample utilities -// (samples/cpp/utils/helpers.h), so mirrored sample code that calls -// `create_cudnn_handle()` works unchanged (RFC 0012 §4.7, §8.3). Unlike the -// upstream helper, this version does not depend on a test framework. +// Mirrors the helper in NVIDIA's sample utilities so mirrored sample code +// compiles unchanged, minus the upstream test-framework dependency. /// @brief RAII deleter that destroys a heap-allocated cuDNN handle. struct CudnnHandleDeleter @@ -132,8 +114,8 @@ struct CudnnHandleDeleter { if(handle != nullptr) { - // A failed destroy at teardown is not recoverable, so we log and - // ignore it (the heap allocation is still freed below). + // A failed destroy at teardown is not recoverable; log and ignore + // (the heap allocation is still freed below). const cudnnStatus_t status = cudnnDestroy(*handle); if(status != CUDNN_STATUS_SUCCESS) { @@ -147,11 +129,8 @@ struct CudnnHandleDeleter /// @brief Create a managed cuDNN handle (mirrors NVIDIA's sample helper). /// -/// The snake_case name intentionally mirrors NVIDIA's helper so mirrored sample -/// code compiles unchanged; the naming check is suppressed accordingly. -/// -/// On a backend create failure the error is logged and an empty pointer is -/// returned, so callers can detect the failure via a null result. +/// snake_case name mirrors NVIDIA's so sample code compiles unchanged. On a +/// backend create failure, logs and returns an empty (null) pointer. inline std::unique_ptr create_cudnn_handle() // NOLINT(readability-identifier-naming) { diff --git a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend.h b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend.h index da897935f0bd..72a1f0fac1d2 100644 --- a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend.h +++ b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend.h @@ -5,30 +5,17 @@ * @file cudnn_frontend.h * @brief Umbrella header for the hipDNN cuDNN-compatibility shim (v9-only). * - * This is the entry point for the cuDNN frontend compatibility shim described in - * RFC 0012 ("cuDNN shim for hipDNN"). It mirrors the filename of NVIDIA's - * upstream `cudnn_frontend.h` so that consumer source can be hipified by - * swapping the include path and (optionally) aliasing the namespace: + * Entry point for the shim. Mirrors the filename of NVIDIA's `cudnn_frontend.h` + * so consumer source can be ported by swapping the include and aliasing the + * namespace: * * @code{.cpp} - * // Textual hipification (RFC §4.2 workflow 1): * #include * namespace cudnn_frontend = hipdnn_frontend::compatibility::cudnn_frontend; * @endcode * - * @note Scope: this shim targets the cuDNN frontend **v9 graph API only** - * (RFC §1, §4.7). It does not reconstruct the v0.x / v8 builder surface. - * - * @note This header must remain includable standalone with no extra `#define`s - * or CMake variables (RFC §4.2). It is installed by the `Development` - * CMake component and is gated, build-side, behind the - * `HIPDNN_ENABLE_CUDNN_COMPATIBILITY` option. - * - * @par Status - * The stub C-API layer — `cudnn.h` (v9-required C-API types, - * handle/stream/error/version entry points, `create_cudnn_handle()`) and the - * version macros — is wired in below. The type/status round-trip mapping, error - * aliasing, and the graph/attribute wrappers land in subsequent tickets. + * Scope: the cuDNN frontend v9 graph API only; the v0.x / v8 builder surface is + * not reconstructed. Must remain includable standalone with no extra defines. */ #pragma once diff --git a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend/graph_helpers.h b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend/graph_helpers.h index 00325e58691b..ad63d9bb79c6 100644 --- a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend/graph_helpers.h +++ b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend/graph_helpers.h @@ -2,29 +2,14 @@ // SPDX-License-Identifier: MIT // // Portions derived from NVIDIA cuDNN frontend (include/cudnn_frontend/graph_helpers.h -// and include/cudnn_frontend_Logging.h), used under the MIT license. The error -// type names, error-code enum, and error/log macro names below intentionally -// mirror cuDNN frontend so hipified v9 consumers compile unchanged. +// and include/cudnn_frontend_Logging.h), used under the MIT license. /** * @file graph_helpers.h * @brief Error-type aliases and error/log macros for the hipDNN cuDNN shim. * - * cuDNN frontend's `error_object` / `error_t` / `error_code_t` (defined upstream - * in `cudnn_frontend/graph_helpers.h`) are structurally identical to - * hipDNN's `Error` / `ErrorCode`, and hipDNN's error-code enum is a - * name-superset of cuDNN's (RFC 0012 §4.6). This header therefore **aliases** - * the hipDNN types into `` with `using` declarations rather than - * wrapping them — `cudnn_frontend::error_t` *is* `hipdnn_frontend::error_t`, so - * an error returned by a wrapped hipDNN call flows out unchanged, with no cast - * between enum families (RFC 0012 §4.3). - * - * It also re-exports the cuDNN frontend error/log macros - * (`CHECK_CUDNN_FRONTEND_ERROR`, `RETURN_CUDNN_FRONTEND_ERROR_IF`, - * `CUDNN_FE_LOG*`) — PyTorch's `AT_CUDNN_FRONTEND_CHECK` expands to them. The - * log macros delegate to hipDNN's frontend logging (off by default, gated by - * `HIPDNN_LOG_LEVEL`); the full cuDNN-FE log env-var bridge (RFC 0012 §4.6) - * lands in a later phase. + * Contains aliases for cuDNN-compatible types from hipDNN, and + * re-exports the cuDNN frontend error/log macros. * * @note Internal-to-shim; pulled in by the umbrella `cudnn_frontend.h`. */ @@ -37,24 +22,12 @@ namespace hipdnn_frontend::compatibility::cudnn_frontend { -// Error types aliased 1:1 (RFC 0012 §4.6). hipDNN's `error_code_t` is a -// name-superset of cuDNN's 16 codes, and `error_object` / `error_t` share -// cuDNN FE's public surface (code, err_msg, get_code/get_message/is_good/ -// is_bad/operator==/!=), so the alias is exact. using hipdnn_frontend::error_code_t; using hipdnn_frontend::error_object; using hipdnn_frontend::error_t; } // namespace hipdnn_frontend::compatibility::cudnn_frontend -// =========================================================================== -// Error/log macros (RFC 0012 §4.6) — re-exported with the cuDNN FE names. -// =========================================================================== -// These are global (un-namespaced), exactly as upstream ships them, because -// consumer code (e.g. PyTorch's AT_CUDNN_FRONTEND_CHECK) uses the bare names. -// `error_t` / `error_code_t` referenced inside expand against whatever -// `cudnn_frontend` alias the consumer has in scope, matching upstream behavior. - // The log macros forward their `<<`-streamed argument straight into hipDNN's // frontend INFO logger. HIPDNN_FE_LOG_INFO gates on HIPDNN_LOG_LEVEL (off by // default) and accepts a stream expression, so chained `<<` arguments work. diff --git a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend/graph_properties.h b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend/graph_properties.h index 4c68e49fd0ef..cb62d6caf845 100644 --- a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend/graph_properties.h +++ b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend/graph_properties.h @@ -2,37 +2,17 @@ // SPDX-License-Identifier: MIT // // Portions derived from NVIDIA cuDNN frontend -// (include/cudnn_frontend/graph_properties.h), used under the MIT license. The -// type name below intentionally mirrors cuDNN frontend so hipified v9 consumers -// compile unchanged. +// (include/cudnn_frontend/graph_properties.h), used under the MIT license. /** * @file graph_properties.h * @brief Graph attribute-type aliases for the hipDNN cuDNN-compatibility shim. * - * cuDNN frontend declares its v9 graph attribute structs (`Tensor_attributes`, - * the per-node `*_attributes`, …) in `cudnn_frontend/graph_properties.h`. This - * header mirrors that filename and brings the matching hipDNN types into - * `::graph` by **aliasing** them with `using` declarations — zero - * overhead, no re-declaration, and no shim-side state (RFC 0012 §4.4.1). - * - * `cudnn_frontend::graph::Tensor_attributes` therefore *is* - * `hipdnn_frontend::graph::TensorAttributes`: a tensor configured through the - * shim flows into a wrapped hipDNN graph with no conversion, and tensor UID - * handling stays entirely on the hipDNN side (`assignUnsetTensorUids()` at - * build). The shim adds no per-tensor identity map or UID allocator - * (RFC 0012 §4.4.1, §5.3). - * - * @par Setter coverage (RFC 0012 §4.4.1) - * hipDNN's `TensorAttributes` publishes the cuDNN-named chained setters the v9 - * graph API uses: `set_dim`, `set_stride`, `set_data_type`, `set_uid`, - * `set_is_virtual`, `set_output`, `set_name` — aliased here 1:1. Pass-by-value - * tensors are expressed through hipDNN's typed `set_value(scalar)` / - * scalar-constructor surface rather than a separate `set_is_pass_by_value(bool)` - * flag. cuDNN FE's `set_is_pass_by_value` and `set_ragged_offset` have no - * hipDNN counterpart yet (no ragged-tensor support); they are intentionally - * out of this alias's scope and are added on the hipDNN side, then surfaced - * here, when their consuming nodes land (RFC 0012 §7.x). + * Brings hipDNN's v9 graph attribute types into `::graph` by aliasing + * them with `using` declarations — zero overhead, no shim-side state. A tensor + * configured through the shim therefore *is* a hipDNN `TensorAttributes` and + * flows into a wrapped hipDNN graph with no conversion (UID handling stays on + * the hipDNN side). * * @note Internal-to-shim; pulled in by the umbrella `cudnn_frontend.h`. */ @@ -44,9 +24,8 @@ namespace hipdnn_frontend::compatibility::cudnn_frontend::graph { -// Tensor attributes aliased 1:1 (RFC 0012 §4.4.1). hipDNN publishes both the -// class name and cuDNN's `Tensor_attributes` spelling (a typedef); both are -// aliased so consumer code using either name resolves through the shim. +// hipDNN publishes cuDNN's `Tensor_attributes` spelling as a typedef; aliasing +// it lets consumer code using that name resolve through the shim. using hipdnn_frontend::graph::Tensor_attributes; } // namespace hipdnn_frontend::compatibility::cudnn_frontend::graph diff --git a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend_utils.h b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend_utils.h index db4d05968c29..ab4edec05f7e 100644 --- a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend_utils.h +++ b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend_utils.h @@ -2,24 +2,17 @@ // SPDX-License-Identifier: MIT // // Portions derived from NVIDIA cuDNN frontend (include/cudnn_frontend_utils.h), -// used under the MIT license. The enum names below intentionally mirror the -// cuDNN frontend FE-namespace enums so hipified v9 consumers compile unchanged. +// used under the MIT license. /** * @file cudnn_frontend_utils.h * @brief FE-namespace enum aliases for the hipDNN cuDNN-compatibility shim. * - * cuDNN frontend uses two parallel enum families (RFC 0012 §4.3): the C-API - * enums from `` and the FE-namespace C++ enums in `namespace - * cudnn_frontend` (`DataType_t`, `PointwiseMode_t`, …). The v9 graph API - * signatures use the **FE-namespace** family; this header makes those names - * available under `` by **aliasing** the matching hipDNN types with - * `using` declarations — zero overhead, no re-declaration, and no numeric cast - * between enum families (RFC 0012 §4.3, §4.3.1). - * - * hipDNN already publishes these enums as cuDNN-named `_t` typedefs (see - * `hipdnn_frontend/Types.hpp`), each a name-superset of its cuDNN counterpart, - * so the alias holds. + * The v9 graph API signatures use cuDNN's FE-namespace enums (`DataType_t`, + * `PointwiseMode_t`, …), not the C-API ones from ``. hipDNN already + * publishes these as cuDNN-named `_t` typedefs (see `Types.hpp`), so this header + * just aliases them into `` with `using` declarations — zero overhead, + * no numeric cast between enum families. * * @note Internal-to-shim; pulled in by the umbrella `cudnn_frontend.h`. */ @@ -31,8 +24,8 @@ namespace hipdnn_frontend::compatibility::cudnn_frontend { -// FE-namespace enums that hipDNN publishes 1:1 (RFC 0012 §4.3). Aliased, not -// re-declared, so `cudnn_frontend::DataType_t` *is* `hipdnn_frontend::DataType_t`. +// FE-namespace enums hipDNN publishes 1:1, aliased so e.g. +// `cudnn_frontend::DataType_t` *is* `hipdnn_frontend::DataType_t`. using hipdnn_frontend::AttentionImplementation_t; using hipdnn_frontend::BehaviorNote_t; using hipdnn_frontend::BuildPlanPolicy_t; @@ -48,17 +41,13 @@ using hipdnn_frontend::ResampleMode_t; // Other cuDNN FE-namespace enums (NumericalNote_t, NormMode_t, RngDistribution_t, // DescriptorType_t, MoeGroupedMatmulMode_t, TensorReordering_t, ReshapeMode_t) -// are intentionally NOT aliased here: hipDNN does not yet publish them, and the -// nodes that use them are out of scope until later phases (RFC 0012 §7.3, §7.4). -// They are added on the hipDNN side and aliased when their node lands. +// are not aliased yet: hipDNN does not publish them and their nodes are out of +// scope. They are aliased when their node lands. } // namespace hipdnn_frontend::compatibility::cudnn_frontend // The `graph` sub-namespace is populated by the `cudnn_frontend/*` headers the -// umbrella pulls in: `graph_properties.h` aliases the attribute types -// (`Tensor_attributes`, …). The `Graph` composition wrapper and the remaining -// per-node `*_attributes` aliases land in later phases (RFC 0012 §7.2). -// Declaring the namespace here keeps this header self-contained even when +// umbrella pulls in; declared empty here so this header is self-contained when // included on its own. namespace hipdnn_frontend::compatibility::cudnn_frontend::graph { diff --git a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend_version.h b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend_version.h index 136ad4073d3d..f456fcf65620 100644 --- a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend_version.h +++ b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_frontend_version.h @@ -6,25 +6,18 @@ /** * @file cudnn_frontend_version.h - * @brief Version macros for the hipDNN cuDNN-compatibility shim. + * @brief cuDNN frontend (FE) version the shim claims source compatibility with. * - * These mirror NVIDIA's `cudnn_frontend_version.h` and declare which cuDNN - * frontend (FE) release this shim claims *source* compatibility with - * (RFC 0012 §4.8). This is the cuDNN **frontend** library version (an - * NVIDIA/cudnn-frontend tag), independent of both hipDNN's own version and the - * cuDNN **runtime** version returned by `cudnnGetVersion()` (see `cudnn.h`). - * - * Consumers such as PyTorch's `MHA.cpp` gate on `CUDNN_FRONTEND_VERSION` - * (e.g. `#if CUDNN_FRONTEND_VERSION <= 11200`), so matching upstream matters. - * - * Pinned to cuDNN FE v1.24.0 (RFC 0012 §2). + * The cuDNN *frontend* library version (a cudnn-frontend tag), distinct from + * hipDNN's version and from the cuDNN *runtime* version in + * `cudnn_runtime_version.h`. Consumers gate on `CUDNN_FRONTEND_VERSION` (e.g. + * PyTorch's `MHA.cpp`), so it must match upstream. Pinned to cuDNN FE v1.24.0. */ #pragma once -// These must remain preprocessor macros (not an enum): consumers such as -// PyTorch's MHA.cpp gate on them in `#if CUDNN_FRONTEND_VERSION <= 11200` -// directives, which an enum cannot satisfy. Suppress modernize-macro-to-enum. +// Must remain preprocessor macros, not an enum: consumers gate on these in `#if` +// directives (e.g. `CUDNN_FRONTEND_VERSION <= 11200`), which an enum cannot do. // NOLINTBEGIN(modernize-macro-to-enum,cppcoreguidelines-macro-to-enum) #define CUDNN_FRONTEND_MAJOR_VERSION 1 #define CUDNN_FRONTEND_MINOR_VERSION 24 diff --git a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_runtime_version.h b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_runtime_version.h index 9121fb2f5d00..023c2daf4877 100644 --- a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_runtime_version.h +++ b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_runtime_version.h @@ -5,29 +5,22 @@ /** * @file cudnn_runtime_version.h - * @brief cuDNN **runtime library** version claimed by the hipDNN shim. + * @brief cuDNN runtime-library version claimed by the shim. * - * NVIDIA's `cudnnGetVersion()` (and the `CUDNN_VERSION` macro) report the cuDNN - * *runtime library* version — e.g. `90500` == 9.5.0 — which is distinct from the - * cuDNN *frontend* version (`CUDNN_FRONTEND_VERSION`, see - * `cudnn_frontend_version.h`). Upstream samples gate on this as a runtime - * version (`cudnnGetVersion() < 90500`, `>= 91400`, ...), so the shim claims a - * cuDNN 9.x runtime version here. `cudnn.h`'s `cudnnGetVersion()` returns - * `CUDNN_VERSION` from this header. + * `cudnnGetVersion()` / `CUDNN_VERSION` report the cuDNN *runtime* version (e.g. + * `90500` == 9.5.0), distinct from the *frontend* version in + * `cudnn_frontend_version.h`. Upstream samples gate on it (`cudnnGetVersion() >= + * 91400`, …), so the shim claims a 9.x runtime here; `cudnn.h`'s + * `cudnnGetVersion()` returns `CUDNN_VERSION`. The claimed 9.22.0 matches what + * cuDNN FE v1.24.0 recommends. * - * The claimed version (9.22.0) matches the runtime the pinned cuDNN frontend - * v1.24.0 release notes recommend (9.22.0 and later). - * - * @note TODO: the exact runtime version claimed (9.22.0) tracks - * the FE v1.24.0 recommendation for now; revisit before attempting a - * PyTorch integration. + * @note TODO: revisit the exact claimed version before a PyTorch integration. */ #pragma once -// These must remain preprocessor macros (not an enum) to mirror NVIDIA's -// `cudnn.h`/`cudnn_version.h`, where consumers gate on `CUDNN_VERSION` in `#if` -// directives that an enum cannot satisfy. Suppress modernize-macro-to-enum. +// Must remain preprocessor macros, not an enum: consumers gate on `CUDNN_VERSION` +// in `#if` directives that an enum cannot satisfy. // NOLINTBEGIN(modernize-macro-to-enum,cppcoreguidelines-macro-to-enum) #define CUDNN_MAJOR 9 #define CUDNN_MINOR 22 diff --git a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_status.h b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_status.h index 6aa178b51b41..5ce851e73942 100644 --- a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_status.h +++ b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/cudnn_status.h @@ -7,11 +7,9 @@ * @file cudnn_status.h * @brief cuDNN library status/return codes for the hipDNN compatibility shim. * - * Split out of `cudnn.h` so that the status enum can be included without pulling - * in the C entry points. `cudnn.h` includes this for its public `cudnnStatus_t` - * type, and `detail/status_translation.h` includes it directly — letting the - * translation header be genuinely self-contained instead of depending on - * `cudnn.h` (which would otherwise create an include cycle). + * Split out of `cudnn.h` so the status enum can be included without the C entry + * points — this lets `detail/status_translation.h` stay self-contained and + * avoids an include cycle. */ #pragma once diff --git a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/detail/status_translation.h b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/detail/status_translation.h index cc6abb704d0d..c49478f256cd 100644 --- a/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/detail/status_translation.h +++ b/projects/hipdnn/frontend/include/hipdnn_compatibility/cudnn/detail/status_translation.h @@ -5,19 +5,13 @@ * @file status_translation.h * @brief Shim-internal translation between cuDNN and hipDNN status codes. * - * Part of the hipDNN cuDNN-compatibility shim (RFC 0012 §4.7). The cuDNN C-API - * stub entry points in `` use these helpers - * to translate return codes on the way through to / back from the hipDNN - * backend. + * The cuDNN C-API stub entry points in `cudnn.h` use these helpers to translate + * return codes to / from the hipDNN backend. The mapping is explicit and named, + * never a numeric cast between the enum families. * - * Identity between the two status enums is established by an explicit named - * mapping, never by a numeric cast between the enum families (RFC 0012 §4.3). - * - * @note Internal header — included by `cudnn.h`, but self-contained: it pulls - * in `cudnn_status.h` for `cudnnStatus_t` and `` for - * `hipdnnStatus_t` directly, rather than depending on `cudnn.h` (which - * would form an include cycle). It lives under `detail/` and is not part - * of the shim's public surface. + * @note Self-contained: pulls in `cudnn_status.h` and `` + * directly rather than `cudnn.h`, to avoid an include cycle. Lives under + * `detail/` and is not part of the shim's public surface. */ #pragma once