diff --git a/projects/hiptensor/CHANGELOG.md b/projects/hiptensor/CHANGELOG.md index b622a25a709d..008af05970e6 100644 --- a/projects/hiptensor/CHANGELOG.md +++ b/projects/hiptensor/CHANGELOG.md @@ -8,6 +8,13 @@ Full documentation for hipTensor is available at [rocm.docs.amd.com/projects/hip * Added support for the following new GPU targets: * gfx11-7: gfx1170, gfx1171. * Added fp64 support to hiptensorPermute for ranks 2-6. +* Added bf16 support to hiptensorPermute for ranks 2-6, with a bf16 or fp32 scalar type. + +### Changed +* Changed `hiptensorPermute` to return `HIPTENSOR_STATUS_NOT_SUPPORTED` when the output tensor's data type differs from the input tensor's. Previously the mismatch passed validation and the call failed with `HIPTENSOR_STATUS_INTERNAL_ERROR` because no kernel matched. + +### Resolved issues +* Fixed `hiptensorElementwiseBinaryExecute` and `hiptensorReduce` treating the alpha, beta, and gamma scalars as 0 when the compute descriptor is `HIPTENSOR_COMPUTE_DESC_16BF`. ## hipTensor 2.5.0 for ROCm 10.2 diff --git a/projects/hiptensor/TESTING.md b/projects/hiptensor/TESTING.md index 9474387917cf..48f65aa47129 100644 --- a/projects/hiptensor/TESTING.md +++ b/projects/hiptensor/TESTING.md @@ -49,16 +49,19 @@ cmake --install --component tests **Purpose:** Validate hardware-independent logic without dispatching a GPU kernel. -Unit tests live in `test/00_unit/` and are built with the `add_hiptensor_unit_test()` CMake function, which links against `hiptensor::hiptensor` and GoogleTest. Most do **not** require a configured GPU to run (the exception is `plan_lifetime_test`, noted below). The test binaries are: +Unit tests live in `test/00_unit/` and are built with the `add_hiptensor_unit_test()` CMake function, which links against `hiptensor::hiptensor` and GoogleTest. Most do **not** launch kernels (the exceptions are `plan_lifetime_test` and `elementwise_broadcast_test`, noted below). The test binaries are: | Binary | What it tests | |---|---| | `logger_test` | Logger singleton identity, callback registration, log-level filtering, file-sink behavior | | `yaml_test` | YAML config parsing and round-trip correctness | | `elementwise_op_test` | CPU-side elementwise operator type traits and identity/scale op correctness | -| `util_test` | Utility helpers (type conversion, index math, stride calculation) | +| `util_test` | Utility helpers (type conversion, scalar reads per compute type, index math, stride calculation) | | `hiptensor_options_test` | CLI/env option parsing (`--hot_runs`, `--cold_runs`, `-y`, `HIPTENSOR_*` env vars) | | `plan_lifetime_test` | Plan object lifetime: using a permutation plan after its descriptor or preference object is destroyed. Despite living in `00_unit/`, it allocates device memory and launches a real permute kernel, so it **requires a GPU** and is registered as a device tier test | +| `batched_contraction_test` | `hiptensorCreateContraction` / `hiptensorCreateContractionTrinary` return `HIPTENSOR_STATUS_NOT_SUPPORTED` for batched modes (a mode shared by both inputs and the output) and accept non-batched ones. No kernel launch | +| `elementwise_mode_test` | Elementwise mode validation: subset and reordered input modes are accepted; a mode missing from the output or repeated within a tensor returns `NOT_SUPPORTED`; a shared mode with mismatched extents returns `INVALID_VALUE`; a permute whose output data type differs from the input's returns `NOT_SUPPORTED`. Allocates small device buffers but launches no kernel | +| `elementwise_broadcast_test` | Elementwise broadcast results for permute, binary, and trinary (inputs carrying a subset of the output modes). Launches kernels, so it **requires a GPU** | > `interface_test` exists in the source tree but is commented out in `test/00_unit/CMakeLists.txt` — it is not currently built. @@ -135,18 +138,19 @@ Integration tests are device tests. They require a physical AMD GPU matching the | Plan cache | `plan_cache_test` | Cache write/read round-trip on device | | Plan lifetime | `plan_lifetime_test` | Permute plan reuse after descriptor/preference destruction (launches a kernel) | | Contraction mode | `contraction_mode_test` | CPU-side mode validation (host test) | -| Elementwise — permute | `rank{2-6}_elementwise_permute_test` | All rank/type combinations | -| Elementwise — binary op | `rank{2-6}_elementwise_binary_op_test` | Binary elementwise ops | -| Elementwise — trinary op | `rank{2-6}_elementwise_trinary_op_test` | Trinary elementwise ops | -| Elementwise — CPU impl | `elementwise_cpu_test` | CPU reference validation (host test; compiles the `elementwise_cpu_impl`, `elementwise_binary_cpu_impl`, and `elementwise_trinary_cpu_impl` sources into one binary) | +| Elementwise — permute | `rank{2-6}_elementwise_permute_test` | FP16/BF16/FP32/FP64; FP16 and BF16 also with an FP32 scalar; ranks 2–6 | +| Elementwise — binary op | `rank{2-6}_elementwise_binary_op_test` | FP16/FP32/FP64; ranks 2–6 (BF16 is accepted by the API but not tested — see Known Gaps) | +| Elementwise — trinary op | `rank{2-6}_elementwise_trinary_op_test` | FP16/FP32/FP64; ranks 2–6 | +| Elementwise — output strides | `elementwise_output_strides_test` | User-supplied (gapped and packed) output strides for permute, binary, and trinary (FP32); launches kernels | +| Elementwise — CPU impl | `elementwise_cpu_test` | CPU reference validation against fixed expected values (host test; compiles the `elementwise_cpu_impl`, `elementwise_binary_cpu_impl`, and `elementwise_trinary_cpu_impl` sources into one binary). Permute: FP32/FP16/BF16 (with a BF16 or FP32 scalar)/FP64 (checked at FP64 precision); binary and trinary: FP32/FP16 | | Reduction | `rank{1-6}_reduction_test` | All rank/type combinations | | Reduction — CPU impl | `reduction_cpu_impl_test` | CPU reference validation (host test) | All device tests compare GPU results against a CPU reference implementation using a per-type tolerance. -**GPU-required tests:** All `bilinear_contraction_test_*`, `scale_contraction_test_*`, `complex_*`, `trinary_*`, `*_unary_ops_test`, `rank*_elementwise_*`, `rank*_reduction_test`, `plan_cache_test`, `plan_lifetime_test`. +**GPU-required tests:** All `bilinear_contraction_test_*`, `scale_contraction_test_*`, `complex_*`, `trinary_*`, `*_unary_ops_test`, `rank*_elementwise_*`, `rank*_reduction_test`, `plan_cache_test`, `plan_lifetime_test`, `elementwise_broadcast_test`, `elementwise_output_strides_test`. -**CPU-only (host) tests:** `contraction_mode_test`, `elementwise_cpu_test`, `reduction_cpu_impl_test`, and the `00_unit/` binaries (`logger_test`, `yaml_test`, `elementwise_op_test`, `util_test`, `hiptensor_options_test`). +**CPU-only (host) tests:** `contraction_mode_test`, `elementwise_cpu_test`, `reduction_cpu_impl_test`, and the `00_unit/` binaries (`logger_test`, `yaml_test`, `elementwise_op_test`, `util_test`, `hiptensor_options_test`). `batched_contraction_test` and `elementwise_mode_test` launch no kernels but create a hipTensor handle and descriptors (`elementwise_mode_test` also allocates small device buffers). **Test tiers (CTest labels):** @@ -268,6 +272,8 @@ llvm-profdata merge -sparse default.profraw -o coverage.profdata llvm-cov report ./bin/ -instr-profile=coverage.profdata ``` +**Coverage test configs:** the `codecoverage` set in `rtest.xml` reads its YAML configs from `/code_coverage_config/`, which the build populates (`copy_code_coverage_config_files` in `cmake/Functions/CopyYamlFolder.cmake`). Contraction and reduction copy from their `configs/code_coverage/` directories; elementwise copies from `configs/emulation/quick/` (permutation, binary_op, trinary_op). `test/02_elementwise/configs/code_coverage/` is not used — edit `emulation/quick` to change elementwise coverage. + **Coverage targets:** | Scope | Target | Notes | @@ -355,7 +361,7 @@ There is no client-side known-bug list (analogous to `known_bugs.yaml` in other | New CPU-side logic (logging, YAML, option parsing, utility math) | Unit test in `test/00_unit/` using GoogleTest | | New or changed GPU operation behavior | Integration test: add a YAML config entry in the appropriate `validation/` tier | | Bug fix | A regression test (unit if CPU-side, integration config entry if GPU-side) that fails before the fix | -| New data type or tensor rank | Config entries across all relevant validation tiers; empty config file for tiers where the test is already covered | +| New data type or tensor rank | Config entries across all relevant validation tiers; empty config file for tiers where the test is already covered. Also add entries to `emulation/quick` **and** `emulation/full` (the FFM tiers do not nest, so `ffm-full` does not inherit `ffm-quick`), to `bench`, and a case in the op's host CPU-reference test (e.g. `elementwise_cpu_test`) where one exists | | Performance-sensitive change | Run `ctest -V -O bench.log` from `/bin/hiptensor/bench` before and after; include the comparison in the PR description. The `scripts/performance/Benchmark*.sh` scripts are an alternative. | | New GPU architecture | Validate on that target; update the Supported Configurations table | @@ -374,6 +380,9 @@ There is no client-side known-bug list (analogous to `known_bugs.yaml` in other | Trinary/unary-ops variants have no bench configs | Low | Low | Not performance-gated | | `clang-format` and static analysis not enforced in CI | Low | Low | Manual enforcement in code review | | No sanitizer (ASAN/TSAN) gate in CI | Low | High if hit | None; ASAN can be built manually | +| Elementwise binary op accepts BF16 but the binary test harness has no BF16 support | Medium | Medium | BF16 scalar reads covered by `util_test`; no end-to-end binary BF16 test | +| Permute `bench` configs exceed device memory on smaller GPUs: the binary aborts with a HIP out-of-memory error partway through (gfx908 and gfx1100 already in the FP32 entries; gfx90a in the FP64 entries at ranks 2 and 4, which are listed last so the other types complete) | Low | Medium | Bench results on those targets are incomplete; run bench on MI300-class GPUs | +| `test/02_elementwise/configs/code_coverage/` is unused (elementwise coverage reads `emulation/quick`) | Low | Low | Documented in the Coverage section | --- diff --git a/projects/hiptensor/docs/api-reference/api-reference.rst b/projects/hiptensor/docs/api-reference/api-reference.rst index e4b6778493c6..a794e343ecdd 100644 --- a/projects/hiptensor/docs/api-reference/api-reference.rst +++ b/projects/hiptensor/docs/api-reference/api-reference.rst @@ -100,6 +100,18 @@ Data Types **** = , where | +------------------------------+ | | | | cf64 / cf64 / cf64 | | | +---------------------+------------------------------+---------------------+---------------------+ +| | f16 / f16 / f16 | gfx9 gfx11 gfx12 | Rank2 - Rank6 | +| Element-wise +------------------------------+ | | +| Permutation | f16 / f16 / f32 | | | +| +------------------------------+ | | +| | bf16 / bf16 / bf16 | | | +| +------------------------------+ | | +| | bf16 / bf16 / f32 | | | +| +------------------------------+ | | +| | f32 / f32 / f32 | | | +| +------------------------------+---------------------+ | +| | f64 / f64 / f64 | gfx90a+ gfx1250 | | ++---------------------+------------------------------+---------------------+---------------------+ | | f16 / f16 / \- | gfx9 gfx11 gfx12 | Rank2 - Rank6 | | Element-wise +------------------------------+ | | | Operations | f16 / f32 / \- | | | diff --git a/projects/hiptensor/docs/install/installation.rst b/projects/hiptensor/docs/install/installation.rst index fa1b408ce266..a6caa60b6b96 100644 --- a/projects/hiptensor/docs/install/installation.rst +++ b/projects/hiptensor/docs/install/installation.rst @@ -341,11 +341,11 @@ Executable name Description ``complex_scale_contraction_test_m4n4k4`` Scale contraction test [D = alpha * (A x B)] with complex single and double precision datatypes of rank 8 ``complex_scale_contraction_test_m5n5k5`` Scale contraction test [D = alpha * (A x B)] with complex single and double precision datatypes of rank 10 ``complex_scale_contraction_test_m6n6k6`` Scale contraction test [D = alpha * (A x B)] with complex single and double precision datatypes of rank 12 -``rank2_elementwise_permute_test`` Permutation test with half, single and double precision datatypes of rank 2 -``rank3_elementwise_permute_test`` Permutation test with half, single and double precision datatypes of rank 3 -``rank4_elementwise_permute_test`` Permutation test with half, single and double precision datatypes of rank 4 -``rank5_elementwise_permute_test`` Permutation test with half, single and double precision datatypes of rank 5 -``rank6_elementwise_permute_test`` Permutation test with half, single and double precision datatypes of rank 6 +``rank2_elementwise_permute_test`` Permutation test with half, bfloat16, single and double precision datatypes of rank 2 +``rank3_elementwise_permute_test`` Permutation test with half, bfloat16, single and double precision datatypes of rank 3 +``rank4_elementwise_permute_test`` Permutation test with half, bfloat16, single and double precision datatypes of rank 4 +``rank5_elementwise_permute_test`` Permutation test with half, bfloat16, single and double precision datatypes of rank 5 +``rank6_elementwise_permute_test`` Permutation test with half, bfloat16, single and double precision datatypes of rank 6 ``rank2_elementwise_binary_op_test`` Element-wise binary operation test with half, single, and double precision datatypes of rank 2 ``rank3_elementwise_binary_op_test`` Element-wise binary operation test with half, single, and double precision datatypes of rank 3 ``rank4_elementwise_binary_op_test`` Element-wise binary operation test with half, single, and double precision datatypes of rank 4 diff --git a/projects/hiptensor/library/src/elementwise/CMakeLists.txt b/projects/hiptensor/library/src/elementwise/CMakeLists.txt index 89d322517747..85a2780e57ec 100644 --- a/projects/hiptensor/library/src/elementwise/CMakeLists.txt +++ b/projects/hiptensor/library/src/elementwise/CMakeLists.txt @@ -67,30 +67,40 @@ set(HIPTENSOR_ELEMENTWISE_SOURCES ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_cpu_reference_rank6_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank2_double_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank2_double_noop_instances.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank2_bhalf_instances.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank2_bhalf_noop_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank2_float_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank2_float_noop_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank2_half_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank2_half_noop_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank3_double_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank3_double_noop_instances.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank3_bhalf_instances.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank3_bhalf_noop_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank3_float_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank3_float_noop_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank3_half_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank3_half_noop_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank4_double_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank4_double_noop_instances.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank4_bhalf_instances.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank4_bhalf_noop_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank4_float_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank4_float_noop_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank4_half_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank4_half_noop_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank5_double_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank5_double_noop_instances.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank5_bhalf_instances.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank5_bhalf_noop_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank5_float_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank5_float_noop_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank5_half_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank5_half_noop_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank6_double_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank6_double_noop_instances.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank6_bhalf_instances.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank6_bhalf_noop_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank6_float_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank6_float_noop_instances.cpp ${CMAKE_CURRENT_SOURCE_DIR}/instances/elementwise_permute_solution_rank6_half_instances.cpp diff --git a/projects/hiptensor/library/src/elementwise/device/hiptensor_elementwise_scale_instances.hpp b/projects/hiptensor/library/src/elementwise/device/hiptensor_elementwise_scale_instances.hpp index ccb872e5337b..fc8f933539f3 100644 --- a/projects/hiptensor/library/src/elementwise/device/hiptensor_elementwise_scale_instances.hpp +++ b/projects/hiptensor/library/src/elementwise/device/hiptensor_elementwise_scale_instances.hpp @@ -319,27 +319,27 @@ namespace ck addInstance<64 , 128 , 32 , 8 , 8 , ck::Sequence<0 , 1> , ck::Sequence<2> , ck::Sequence<2>>(opPtrs); addInstance<64 , 128 , 32 , 8 , 8 , ck::Sequence<0 , 1> , ck::Sequence<1> , ck::Sequence<1>>(opPtrs); - } else if constexpr(std::is_same_v> && NumDim == 2) { + } else if constexpr((std::is_same_v> || std::is_same_v>) && NumDim == 2) { addInstance<256 , 64 , 64 , 4 , 4 , ck::Sequence<0 , 1> , ck::Sequence<4> , ck::Sequence<4>>(opPtrs); - // the following instances are the safety net to double and rank2 + // the following instances are the safety net to double/bhalf and rank2 addInstance<256 , 64 , 64 , 4 , 4 , ck::Sequence<0 , 1> , ck::Sequence<2> , ck::Sequence<2>>(opPtrs); addInstance<256 , 64 , 64 , 4 , 4 , ck::Sequence<0 , 1> , ck::Sequence<1> , ck::Sequence<1>>(opPtrs); - } else if constexpr(std::is_same_v> && NumDim == 3) { + } else if constexpr((std::is_same_v> || std::is_same_v>) && NumDim == 3) { addInstance<256 , 64 , 64 , 4 , 4 , ck::Sequence<0 , 1> , ck::Sequence<4> , ck::Sequence<4>>(opPtrs); - // the following instances are the safety net to double and rank3 + // the following instances are the safety net to double/bhalf and rank3 addInstance<256 , 64 , 64 , 4 , 4 , ck::Sequence<0 , 1> , ck::Sequence<2> , ck::Sequence<2>>(opPtrs); addInstance<256 , 64 , 64 , 4 , 4 , ck::Sequence<0 , 1> , ck::Sequence<1> , ck::Sequence<1>>(opPtrs); - } else if constexpr(std::is_same_v> && NumDim == 4) { + } else if constexpr((std::is_same_v> || std::is_same_v>) && NumDim == 4) { addInstance<256 , 64 , 64 , 4 , 4 , ck::Sequence<0 , 1> , ck::Sequence<4> , ck::Sequence<4>>(opPtrs); - // the following instances are the safety net to double and rank4 + // the following instances are the safety net to double/bhalf and rank4 addInstance<256 , 64 , 64 , 4 , 4 , ck::Sequence<0 , 1> , ck::Sequence<2> , ck::Sequence<2>>(opPtrs); addInstance<256 , 64 , 64 , 4 , 4 , ck::Sequence<0 , 1> , ck::Sequence<1> , ck::Sequence<1>>(opPtrs); diff --git a/projects/hiptensor/library/src/elementwise/device/instance_params.cpp b/projects/hiptensor/library/src/elementwise/device/instance_params.cpp index 705acab2e75b..1685f9731acf 100644 --- a/projects/hiptensor/library/src/elementwise/device/instance_params.cpp +++ b/projects/hiptensor/library/src/elementwise/device/instance_params.cpp @@ -52,7 +52,7 @@ namespace ck::tensor_operation::device::instance } else if (typeIn[0] == HIPTENSOR_R_32F) { hashCodes.push_back(hiptensor::Hash{}( DeviceElementwiseParams::Gen(typeIn , typeOut , scale , numDim , {256 , 64 , 64 , 4 , 4 , {0 , 1} , {2} , {2}}))); hashCodes.push_back(hiptensor::Hash{}( DeviceElementwiseParams::Gen(typeIn , typeOut , scale , numDim , {256 , 64 , 64 , 4 , 4 , {0 , 1} , {1} , {1}}))); - } else if (typeIn[0] == HIPTENSOR_R_64F) { + } else if (typeIn[0] == HIPTENSOR_R_64F || typeIn[0] == HIPTENSOR_R_16BF) { hashCodes.push_back(hiptensor::Hash{}( DeviceElementwiseParams::Gen(typeIn , typeOut , scale , numDim , {256 , 64 , 64 , 4 , 4 , {0 , 1} , {2} , {2}}))); hashCodes.push_back(hiptensor::Hash{}( DeviceElementwiseParams::Gen(typeIn , typeOut , scale , numDim , {256 , 64 , 64 , 4 , 4 , {0 , 1} , {1} , {1}}))); } @@ -63,7 +63,7 @@ namespace ck::tensor_operation::device::instance } else if (typeIn[0] == HIPTENSOR_R_32F) { hashCodes.push_back(hiptensor::Hash{}( DeviceElementwiseParams::Gen(typeIn , typeOut , scale , numDim , {256 , 64 , 64 , 4 , 4 , {0 , 1} , {2} , {2}}))); hashCodes.push_back(hiptensor::Hash{}( DeviceElementwiseParams::Gen(typeIn , typeOut , scale , numDim , {256 , 64 , 64 , 4 , 4 , {0 , 1} , {1} , {1}}))); - } else if (typeIn[0] == HIPTENSOR_R_64F) { + } else if (typeIn[0] == HIPTENSOR_R_64F || typeIn[0] == HIPTENSOR_R_16BF) { hashCodes.push_back(hiptensor::Hash{}( DeviceElementwiseParams::Gen(typeIn , typeOut , scale , numDim , {256 , 64 , 64 , 4 , 4 , {0 , 1} , {2} , {2}}))); hashCodes.push_back(hiptensor::Hash{}( DeviceElementwiseParams::Gen(typeIn , typeOut , scale , numDim , {256 , 64 , 64 , 4 , 4 , {0 , 1} , {1} , {1}}))); } @@ -74,7 +74,7 @@ namespace ck::tensor_operation::device::instance } else if (typeIn[0] == HIPTENSOR_R_32F) { hashCodes.push_back(hiptensor::Hash{}( DeviceElementwiseParams::Gen(typeIn , typeOut , scale , numDim , {256 , 64 , 64 , 4 , 4 , {0 , 1} , {2} , {2}}))); hashCodes.push_back(hiptensor::Hash{}( DeviceElementwiseParams::Gen(typeIn , typeOut , scale , numDim , {256 , 64 , 64 , 4 , 4 , {0 , 1} , {1} , {1}}))); - } else if (typeIn[0] == HIPTENSOR_R_64F) { + } else if (typeIn[0] == HIPTENSOR_R_64F || typeIn[0] == HIPTENSOR_R_16BF) { hashCodes.push_back(hiptensor::Hash{}( DeviceElementwiseParams::Gen(typeIn , typeOut , scale , numDim , {256 , 64 , 64 , 4 , 4 , {0 , 1} , {2} , {2}}))); hashCodes.push_back(hiptensor::Hash{}( DeviceElementwiseParams::Gen(typeIn , typeOut , scale , numDim , {256 , 64 , 64 , 4 , 4 , {0 , 1} , {1} , {1}}))); } diff --git a/projects/hiptensor/library/src/elementwise/elementwise_solution_instances.cpp b/projects/hiptensor/library/src/elementwise/elementwise_solution_instances.cpp index 7c79c4395646..532a54464b45 100644 --- a/projects/hiptensor/library/src/elementwise/elementwise_solution_instances.cpp +++ b/projects/hiptensor/library/src/elementwise/elementwise_solution_instances.cpp @@ -62,18 +62,24 @@ namespace hiptensor ElementwiseTrinarySolution6DHalfInstances(); ElementwisePermuteSolution4DDoubleInstances(); ElementwisePermuteSolution4DDoubleNoopInstances(); + ElementwisePermuteSolution4DBhalfInstances(); + ElementwisePermuteSolution4DBhalfNoopInstances(); ElementwisePermuteSolution4DFloatInstances(); ElementwisePermuteSolution4DFloatNoopInstances(); ElementwisePermuteSolution4DHalfInstances(); ElementwisePermuteSolution4DHalfNoopInstances(); ElementwisePermuteSolution5DDoubleInstances(); ElementwisePermuteSolution5DDoubleNoopInstances(); + ElementwisePermuteSolution5DBhalfInstances(); + ElementwisePermuteSolution5DBhalfNoopInstances(); ElementwisePermuteSolution5DFloatInstances(); ElementwisePermuteSolution5DFloatNoopInstances(); ElementwisePermuteSolution5DHalfInstances(); ElementwisePermuteSolution5DHalfNoopInstances(); ElementwisePermuteSolution6DDoubleInstances(); ElementwisePermuteSolution6DDoubleNoopInstances(); + ElementwisePermuteSolution6DBhalfInstances(); + ElementwisePermuteSolution6DBhalfNoopInstances(); ElementwisePermuteSolution6DFloatInstances(); ElementwisePermuteSolution6DFloatNoopInstances(); ElementwisePermuteSolution6DHalfInstances(); @@ -84,12 +90,16 @@ namespace hiptensor ElementwiseBinarySolution3DHalfInstances(); ElementwisePermuteSolution2DDoubleInstances(); ElementwisePermuteSolution2DDoubleNoopInstances(); + ElementwisePermuteSolution2DBhalfInstances(); + ElementwisePermuteSolution2DBhalfNoopInstances(); ElementwisePermuteSolution2DFloatInstances(); ElementwisePermuteSolution2DFloatNoopInstances(); ElementwisePermuteSolution2DHalfInstances(); ElementwisePermuteSolution2DHalfNoopInstances(); ElementwisePermuteSolution3DDoubleInstances(); ElementwisePermuteSolution3DDoubleNoopInstances(); + ElementwisePermuteSolution3DBhalfInstances(); + ElementwisePermuteSolution3DBhalfNoopInstances(); ElementwisePermuteSolution3DFloatInstances(); ElementwisePermuteSolution3DFloatNoopInstances(); ElementwisePermuteSolution3DHalfInstances(); diff --git a/projects/hiptensor/library/src/elementwise/elementwise_solution_instances.hpp b/projects/hiptensor/library/src/elementwise/elementwise_solution_instances.hpp index c38cf0704469..c04c19d56e89 100644 --- a/projects/hiptensor/library/src/elementwise/elementwise_solution_instances.hpp +++ b/projects/hiptensor/library/src/elementwise/elementwise_solution_instances.hpp @@ -76,30 +76,40 @@ namespace hiptensor void ElementwiseTrinarySolution6DHalfInstances(); void ElementwisePermuteSolution2DDoubleInstances(); void ElementwisePermuteSolution2DDoubleNoopInstances(); + void ElementwisePermuteSolution2DBhalfInstances(); + void ElementwisePermuteSolution2DBhalfNoopInstances(); void ElementwisePermuteSolution2DFloatInstances(); void ElementwisePermuteSolution2DFloatNoopInstances(); void ElementwisePermuteSolution2DHalfInstances(); void ElementwisePermuteSolution2DHalfNoopInstances(); void ElementwisePermuteSolution3DDoubleInstances(); void ElementwisePermuteSolution3DDoubleNoopInstances(); + void ElementwisePermuteSolution3DBhalfInstances(); + void ElementwisePermuteSolution3DBhalfNoopInstances(); void ElementwisePermuteSolution3DFloatInstances(); void ElementwisePermuteSolution3DFloatNoopInstances(); void ElementwisePermuteSolution3DHalfInstances(); void ElementwisePermuteSolution3DHalfNoopInstances(); void ElementwisePermuteSolution4DDoubleInstances(); void ElementwisePermuteSolution4DDoubleNoopInstances(); + void ElementwisePermuteSolution4DBhalfInstances(); + void ElementwisePermuteSolution4DBhalfNoopInstances(); void ElementwisePermuteSolution4DFloatInstances(); void ElementwisePermuteSolution4DFloatNoopInstances(); void ElementwisePermuteSolution4DHalfInstances(); void ElementwisePermuteSolution4DHalfNoopInstances(); void ElementwisePermuteSolution5DDoubleInstances(); void ElementwisePermuteSolution5DDoubleNoopInstances(); + void ElementwisePermuteSolution5DBhalfInstances(); + void ElementwisePermuteSolution5DBhalfNoopInstances(); void ElementwisePermuteSolution5DFloatInstances(); void ElementwisePermuteSolution5DFloatNoopInstances(); void ElementwisePermuteSolution5DHalfInstances(); void ElementwisePermuteSolution5DHalfNoopInstances(); void ElementwisePermuteSolution6DDoubleInstances(); void ElementwisePermuteSolution6DDoubleNoopInstances(); + void ElementwisePermuteSolution6DBhalfInstances(); + void ElementwisePermuteSolution6DBhalfNoopInstances(); void ElementwisePermuteSolution6DFloatInstances(); void ElementwisePermuteSolution6DFloatNoopInstances(); void ElementwisePermuteSolution6DHalfInstances(); diff --git a/projects/hiptensor/library/src/elementwise/hiptensor_elementwise_permute.cpp b/projects/hiptensor/library/src/elementwise/hiptensor_elementwise_permute.cpp index 1c2ce68bbda2..50410a2f157d 100644 --- a/projects/hiptensor/library/src/elementwise/hiptensor_elementwise_permute.cpp +++ b/projects/hiptensor/library/src/elementwise/hiptensor_elementwise_permute.cpp @@ -88,17 +88,20 @@ hiptensorStatus_t hiptensorPermute(const hiptensorHandle_t handle, return checkResult; } - constexpr std::array, 4> validDataTypes - = {{// typeA, typeC, typeScalar + constexpr std::array, 6> validDataTypes + = {{// typeA, typeScalar {HIPTENSOR_R_16F, HIPTENSOR_R_16F}, {HIPTENSOR_R_16F, HIPTENSOR_R_32F}, + {HIPTENSOR_R_16BF, HIPTENSOR_R_16BF}, + {HIPTENSOR_R_16BF, HIPTENSOR_R_32F}, {HIPTENSOR_R_32F, HIPTENSOR_R_32F}, {HIPTENSOR_R_64F, HIPTENSOR_R_64F}}}; std::array inputTensorTypes = {descA->mType, typeScalar}; - if(std::none_of(validDataTypes.cbegin(), - validDataTypes.cend(), - [&inputTensorTypes](auto&& types) { return types == inputTensorTypes; })) + if(descA->mType != descB->mType + || std::none_of(validDataTypes.cbegin(), + validDataTypes.cend(), + [&inputTensorTypes](auto&& types) { return types == inputTensorTypes; })) { auto errorCode = HIPTENSOR_STATUS_NOT_SUPPORTED; snprintf(msg, diff --git a/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank2_bhalf_instances.cpp b/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank2_bhalf_instances.cpp new file mode 100644 index 000000000000..9b4f374425bb --- /dev/null +++ b/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank2_bhalf_instances.cpp @@ -0,0 +1,46 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (C) 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. + * + *******************************************************************************/ + +#include "../elementwise_solution.hpp" +#include "../elementwise_solution_instances.hpp" + +// Ensure access to +#include "../device/hiptensor_elementwise_scale_instances.hpp" +#include "hiptensor_ck_types.hpp" +#include + +namespace hiptensor +{ + void ElementwiseSolutionInstances::ElementwisePermuteSolution2DBhalfInstances() + { + // Register all the solutions exactly once + // 2d Permutation + registerSolutions(enumerateElementwiseSolutions, + ck::Tuple, + CkPermutationUnaryCombinedOp, + 2>()); + } +} // namespace hiptensor diff --git a/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank2_bhalf_noop_instances.cpp b/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank2_bhalf_noop_instances.cpp new file mode 100644 index 000000000000..81576b2df445 --- /dev/null +++ b/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank2_bhalf_noop_instances.cpp @@ -0,0 +1,45 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (C) 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. + * + *******************************************************************************/ + +#include "../elementwise_solution.hpp" +#include "../elementwise_solution_instances.hpp" + +// Ensure access to +#include "../device/hiptensor_elementwise_scale_instances.hpp" +#include "hiptensor_ck_types.hpp" + +namespace hiptensor +{ + void ElementwiseSolutionInstances::ElementwisePermuteSolution2DBhalfNoopInstances() + { + // Register all the solutions exactly once + // 2d Permutation + registerSolutions(enumerateElementwiseSolutions, + ck::Tuple, + CkPermutationPassThroughCombinedOp, + 2>()); + } +} // namespace hiptensor diff --git a/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank3_bhalf_instances.cpp b/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank3_bhalf_instances.cpp new file mode 100644 index 000000000000..7f50ef7f9e11 --- /dev/null +++ b/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank3_bhalf_instances.cpp @@ -0,0 +1,46 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (C) 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. + * + *******************************************************************************/ + +#include "../elementwise_solution.hpp" +#include "../elementwise_solution_instances.hpp" + +// Ensure access to +#include "../device/hiptensor_elementwise_scale_instances.hpp" +#include "hiptensor_ck_types.hpp" +#include + +namespace hiptensor +{ + void ElementwiseSolutionInstances::ElementwisePermuteSolution3DBhalfInstances() + { + // Register all the solutions exactly once + // 3d Permutation + registerSolutions(enumerateElementwiseSolutions, + ck::Tuple, + CkPermutationUnaryCombinedOp, + 3>()); + } +} // namespace hiptensor diff --git a/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank3_bhalf_noop_instances.cpp b/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank3_bhalf_noop_instances.cpp new file mode 100644 index 000000000000..30ec46ed8f57 --- /dev/null +++ b/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank3_bhalf_noop_instances.cpp @@ -0,0 +1,45 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (C) 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. + * + *******************************************************************************/ + +#include "../elementwise_solution.hpp" +#include "../elementwise_solution_instances.hpp" + +// Ensure access to +#include "../device/hiptensor_elementwise_scale_instances.hpp" +#include "hiptensor_ck_types.hpp" + +namespace hiptensor +{ + void ElementwiseSolutionInstances::ElementwisePermuteSolution3DBhalfNoopInstances() + { + // Register all the solutions exactly once + // 3d Permutation + registerSolutions(enumerateElementwiseSolutions, + ck::Tuple, + CkPermutationPassThroughCombinedOp, + 3>()); + } +} // namespace hiptensor diff --git a/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank4_bhalf_instances.cpp b/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank4_bhalf_instances.cpp new file mode 100644 index 000000000000..0826174a9760 --- /dev/null +++ b/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank4_bhalf_instances.cpp @@ -0,0 +1,46 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (C) 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. + * + *******************************************************************************/ + +#include "../elementwise_solution.hpp" +#include "../elementwise_solution_instances.hpp" + +// Ensure access to +#include "../device/hiptensor_elementwise_scale_instances.hpp" +#include "hiptensor_ck_types.hpp" +#include + +namespace hiptensor +{ + void ElementwiseSolutionInstances::ElementwisePermuteSolution4DBhalfInstances() + { + // Register all the solutions exactly once + // 4d Permutation + registerSolutions(enumerateElementwiseSolutions, + ck::Tuple, + CkPermutationUnaryCombinedOp, + 4>()); + } +} // namespace hiptensor diff --git a/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank4_bhalf_noop_instances.cpp b/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank4_bhalf_noop_instances.cpp new file mode 100644 index 000000000000..07c7c84e1884 --- /dev/null +++ b/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank4_bhalf_noop_instances.cpp @@ -0,0 +1,45 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (C) 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. + * + *******************************************************************************/ + +#include "../elementwise_solution.hpp" +#include "../elementwise_solution_instances.hpp" + +// Ensure access to +#include "../device/hiptensor_elementwise_scale_instances.hpp" +#include "hiptensor_ck_types.hpp" + +namespace hiptensor +{ + void ElementwiseSolutionInstances::ElementwisePermuteSolution4DBhalfNoopInstances() + { + // Register all the solutions exactly once + // 4d Permutation + registerSolutions(enumerateElementwiseSolutions, + ck::Tuple, + CkPermutationPassThroughCombinedOp, + 4>()); + } +} // namespace hiptensor diff --git a/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank5_bhalf_instances.cpp b/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank5_bhalf_instances.cpp new file mode 100644 index 000000000000..6652de14ccbe --- /dev/null +++ b/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank5_bhalf_instances.cpp @@ -0,0 +1,46 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (C) 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. + * + *******************************************************************************/ + +#include "../elementwise_solution.hpp" +#include "../elementwise_solution_instances.hpp" + +// Ensure access to +#include "../device/hiptensor_elementwise_scale_instances.hpp" +#include "hiptensor_ck_types.hpp" +#include + +namespace hiptensor +{ + void ElementwiseSolutionInstances::ElementwisePermuteSolution5DBhalfInstances() + { + // Register all the solutions exactly once + // 5d Permutation + registerSolutions(enumerateElementwiseSolutions, + ck::Tuple, + CkPermutationUnaryCombinedOp, + 5>()); + } +} // namespace hiptensor diff --git a/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank5_bhalf_noop_instances.cpp b/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank5_bhalf_noop_instances.cpp new file mode 100644 index 000000000000..fe76dd81bc59 --- /dev/null +++ b/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank5_bhalf_noop_instances.cpp @@ -0,0 +1,45 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (C) 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. + * + *******************************************************************************/ + +#include "../elementwise_solution.hpp" +#include "../elementwise_solution_instances.hpp" + +// Ensure access to +#include "../device/hiptensor_elementwise_scale_instances.hpp" +#include "hiptensor_ck_types.hpp" + +namespace hiptensor +{ + void ElementwiseSolutionInstances::ElementwisePermuteSolution5DBhalfNoopInstances() + { + // Register all the solutions exactly once + // 5d Permutation + registerSolutions(enumerateElementwiseSolutions, + ck::Tuple, + CkPermutationPassThroughCombinedOp, + 5>()); + } +} // namespace hiptensor diff --git a/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank6_bhalf_instances.cpp b/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank6_bhalf_instances.cpp new file mode 100644 index 000000000000..17faebb86e09 --- /dev/null +++ b/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank6_bhalf_instances.cpp @@ -0,0 +1,46 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (C) 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. + * + *******************************************************************************/ + +#include "../elementwise_solution.hpp" +#include "../elementwise_solution_instances.hpp" + +// Ensure access to +#include "../device/hiptensor_elementwise_scale_instances.hpp" +#include "hiptensor_ck_types.hpp" +#include + +namespace hiptensor +{ + void ElementwiseSolutionInstances::ElementwisePermuteSolution6DBhalfInstances() + { + // Register all the solutions exactly once + // 6d Permutation + registerSolutions(enumerateElementwiseSolutions, + ck::Tuple, + CkPermutationUnaryCombinedOp, + 6>()); + } +} // namespace hiptensor diff --git a/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank6_bhalf_noop_instances.cpp b/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank6_bhalf_noop_instances.cpp new file mode 100644 index 000000000000..6d9c18511678 --- /dev/null +++ b/projects/hiptensor/library/src/elementwise/instances/elementwise_permute_solution_rank6_bhalf_noop_instances.cpp @@ -0,0 +1,45 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (C) 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. + * + *******************************************************************************/ + +#include "../elementwise_solution.hpp" +#include "../elementwise_solution_instances.hpp" + +// Ensure access to +#include "../device/hiptensor_elementwise_scale_instances.hpp" +#include "hiptensor_ck_types.hpp" + +namespace hiptensor +{ + void ElementwiseSolutionInstances::ElementwisePermuteSolution6DBhalfNoopInstances() + { + // Register all the solutions exactly once + // 6d Permutation + registerSolutions(enumerateElementwiseSolutions, + ck::Tuple, + CkPermutationPassThroughCombinedOp, + 6>()); + } +} // namespace hiptensor diff --git a/projects/hiptensor/library/src/include/data_types_impl.hpp b/projects/hiptensor/library/src/include/data_types_impl.hpp index 54e9bebfeae3..1f9d436fbd4b 100644 --- a/projects/hiptensor/library/src/include/data_types_impl.hpp +++ b/projects/hiptensor/library/src/include/data_types_impl.hpp @@ -195,7 +195,11 @@ namespace hiptensor template T readVal(void const* value, hiptensorComputeDescriptor_t id) { - if(id == HIPTENSOR_COMPUTE_DESC_16F) + if(id == HIPTENSOR_COMPUTE_DESC_16BF) + { + return static_cast(static_cast(*(hip_bfloat16*)value)); + } + else if(id == HIPTENSOR_COMPUTE_DESC_16F) { return static_cast(*(_Float16*)value); } diff --git a/projects/hiptensor/test/00_unit/elementwise_mode_test.cpp b/projects/hiptensor/test/00_unit/elementwise_mode_test.cpp index eac2d5c9cb2a..79007a07141b 100644 --- a/projects/hiptensor/test/00_unit/elementwise_mode_test.cpp +++ b/projects/hiptensor/test/00_unit/elementwise_mode_test.cpp @@ -165,6 +165,51 @@ TEST_F(ElementwiseModeTest, PermutationWithReorderedModesIsSupported) hiptensorDestroyOperationDescriptor(opDesc); } +// Permute does not convert types, so an output whose data type differs from the input's has no +// matching kernel and must be rejected up front. +TEST_F(ElementwiseModeTest, PermutationWithMismatchedOutputTypeReturnsNotSupported) +{ + int64_t lengths[2] = {2, 3}; + hiptensorTensorDescriptor_t bf16Desc{}; + ASSERT_EQ(hiptensorCreateTensorDescriptor( + mHandle, &bf16Desc, 2, lengths, nullptr, HIPTENSOR_R_16BF, alignment), + HIPTENSOR_STATUS_SUCCESS); + + hiptensorOperationDescriptor_t opDesc{}; + ASSERT_EQ(hiptensorCreatePermutation(mHandle, + &opDesc, + bf16Desc, + mFullModes, + HIPTENSOR_OP_IDENTITY, + mFull, + mFullModes, + HIPTENSOR_COMPUTE_DESC_32F), + HIPTENSOR_STATUS_SUCCESS); + + hiptensorPlanPreference_t pref{}; + ASSERT_EQ(hiptensorCreatePlanPreference( + mHandle, &pref, HIPTENSOR_ALGO_DEFAULT, HIPTENSOR_JIT_MODE_NONE), + HIPTENSOR_STATUS_SUCCESS); + hiptensorPlan_t plan{}; + ASSERT_EQ(hiptensorCreatePlan(mHandle, &plan, opDesc, pref, 0), HIPTENSOR_STATUS_SUCCESS); + + void* a = nullptr; + void* b = nullptr; + ASSERT_EQ(hipMalloc(&a, 2 * 3 * sizeof(uint16_t)), hipSuccess); + ASSERT_EQ(hipMalloc(&b, 2 * 3 * sizeof(float)), hipSuccess); + + float alpha = 2.0f; + EXPECT_EQ(hiptensorPermute(mHandle, plan, &alpha, a, b, nullptr), + HIPTENSOR_STATUS_NOT_SUPPORTED); + + EXPECT_EQ(hipFree(b), hipSuccess); + EXPECT_EQ(hipFree(a), hipSuccess); + hiptensorDestroyPlan(plan); + hiptensorDestroyPlanPreference(pref); + hiptensorDestroyOperationDescriptor(opDesc); + hiptensorDestroyTensorDescriptor(bf16Desc); +} + TEST_F(ElementwiseModeTest, BinaryWithSubsetModesIsSupported) { hiptensorOperationDescriptor_t opDesc{}; diff --git a/projects/hiptensor/test/00_unit/util_test.cpp b/projects/hiptensor/test/00_unit/util_test.cpp index 36d6e856fbf8..73567a4041bf 100644 --- a/projects/hiptensor/test/00_unit/util_test.cpp +++ b/projects/hiptensor/test/00_unit/util_test.cpp @@ -148,6 +148,19 @@ TEST(convertToHipTensorDataTypeTypeToStringTest, UtilTest) EXPECT_EQ(hiptensor::convertToHipTensorDataType(HIPTENSOR_COMPUTE_DESC_C64F), HIPTENSOR_C_64F); } +TEST(readValComputeDescTest, UtilTest) +{ + hip_bfloat16 bf16Val(2.5f); + _Float16 f16Val = 2.5; + float f32Val = 2.5f; + double f64Val = 2.5; + EXPECT_EQ(hiptensor::readVal(&bf16Val, HIPTENSOR_COMPUTE_DESC_16BF), 2.5f); + EXPECT_EQ(hiptensor::readVal(&bf16Val, HIPTENSOR_COMPUTE_DESC_16BF), 2.5); + EXPECT_EQ(hiptensor::readVal(&f16Val, HIPTENSOR_COMPUTE_DESC_16F), 2.5); + EXPECT_EQ(hiptensor::readVal(&f32Val, HIPTENSOR_COMPUTE_DESC_32F), 2.5); + EXPECT_EQ(hiptensor::readVal(&f64Val, HIPTENSOR_COMPUTE_DESC_64F), 2.5); +} + TEST(algoTypeToStringTest, UtilTest) { EXPECT_STREQ(hiptensor::algoTypeToString(HIPTENSOR_ALGO_ACTOR_CRITIC).c_str(), diff --git a/projects/hiptensor/test/02_elementwise/configs/bench/rank2_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/bench/rank2_test_params.yaml index e0035003ad3f..b262d4687b3e 100644 --- a/projects/hiptensor/test/02_elementwise/configs/bench/rank2_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/bench/rank2_test_params.yaml @@ -4,6 +4,9 @@ Tensor Data Types: - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_64F, HIPTENSOR_R_64F] Alphas: - 0.0 - 1.2 diff --git a/projects/hiptensor/test/02_elementwise/configs/bench/rank3_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/bench/rank3_test_params.yaml index ce4a5885521a..c077aee944f0 100644 --- a/projects/hiptensor/test/02_elementwise/configs/bench/rank3_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/bench/rank3_test_params.yaml @@ -4,6 +4,9 @@ Tensor Data Types: - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_64F, HIPTENSOR_R_64F] Alphas: - 0.0 - 1.2 diff --git a/projects/hiptensor/test/02_elementwise/configs/bench/rank4_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/bench/rank4_test_params.yaml index c0c789c5b11d..59b009a7c2c8 100644 --- a/projects/hiptensor/test/02_elementwise/configs/bench/rank4_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/bench/rank4_test_params.yaml @@ -4,6 +4,9 @@ Tensor Data Types: - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_64F, HIPTENSOR_R_64F] Alphas: - 0.0 - 1.2 diff --git a/projects/hiptensor/test/02_elementwise/configs/bench/rank5_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/bench/rank5_test_params.yaml index 735f18db2ee6..ada0485527f5 100644 --- a/projects/hiptensor/test/02_elementwise/configs/bench/rank5_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/bench/rank5_test_params.yaml @@ -4,6 +4,9 @@ Tensor Data Types: - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_64F, HIPTENSOR_R_64F] Alphas: - 0.0 - 1.2 diff --git a/projects/hiptensor/test/02_elementwise/configs/bench/rank6_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/bench/rank6_test_params.yaml index 7ea89aa47ceb..3b9f5f065b90 100644 --- a/projects/hiptensor/test/02_elementwise/configs/bench/rank6_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/bench/rank6_test_params.yaml @@ -4,6 +4,9 @@ Tensor Data Types: - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_64F, HIPTENSOR_R_64F] Alphas: - 0.0 - 1.2 diff --git a/projects/hiptensor/test/02_elementwise/configs/emulation/full/permutation/rank2_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/emulation/full/permutation/rank2_test_params.yaml index 09901a30d315..6a8bc2328003 100644 --- a/projects/hiptensor/test/02_elementwise/configs/emulation/full/permutation/rank2_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/emulation/full/permutation/rank2_test_params.yaml @@ -1,9 +1,12 @@ --- Log Level: [ HIPTENSOR_LOG_LEVEL_ERROR, HIPTENSOR_LOG_LEVEL_PERF_TRACE ] Tensor Data Types: + - [ HIPTENSOR_R_64F, HIPTENSOR_R_64F] - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_32F] Alphas: - 2.3 Lengths: diff --git a/projects/hiptensor/test/02_elementwise/configs/emulation/full/permutation/rank3_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/emulation/full/permutation/rank3_test_params.yaml index 4af856a20e0e..f1da8c78b9a7 100644 --- a/projects/hiptensor/test/02_elementwise/configs/emulation/full/permutation/rank3_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/emulation/full/permutation/rank3_test_params.yaml @@ -1,9 +1,12 @@ --- Log Level: [ HIPTENSOR_LOG_LEVEL_ERROR, HIPTENSOR_LOG_LEVEL_PERF_TRACE ] Tensor Data Types: + - [ HIPTENSOR_R_64F, HIPTENSOR_R_64F] - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_32F] Alphas: - 2.3 Lengths: diff --git a/projects/hiptensor/test/02_elementwise/configs/emulation/full/permutation/rank4_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/emulation/full/permutation/rank4_test_params.yaml index 7a232eddd144..02265fad427f 100644 --- a/projects/hiptensor/test/02_elementwise/configs/emulation/full/permutation/rank4_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/emulation/full/permutation/rank4_test_params.yaml @@ -1,9 +1,12 @@ --- Log Level: [ HIPTENSOR_LOG_LEVEL_ERROR, HIPTENSOR_LOG_LEVEL_PERF_TRACE ] Tensor Data Types: + - [ HIPTENSOR_R_64F, HIPTENSOR_R_64F] - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_32F] Alphas: - 2.3 Lengths: diff --git a/projects/hiptensor/test/02_elementwise/configs/emulation/full/permutation/rank5_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/emulation/full/permutation/rank5_test_params.yaml index a582dd73f60f..f0d70b7d1378 100644 --- a/projects/hiptensor/test/02_elementwise/configs/emulation/full/permutation/rank5_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/emulation/full/permutation/rank5_test_params.yaml @@ -1,9 +1,12 @@ --- Log Level: [ HIPTENSOR_LOG_LEVEL_ERROR, HIPTENSOR_LOG_LEVEL_PERF_TRACE ] Tensor Data Types: + - [ HIPTENSOR_R_64F, HIPTENSOR_R_64F] - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_32F] Alphas: - 2.3 Lengths: diff --git a/projects/hiptensor/test/02_elementwise/configs/emulation/full/permutation/rank6_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/emulation/full/permutation/rank6_test_params.yaml index c77634b2605d..9fc015e94054 100644 --- a/projects/hiptensor/test/02_elementwise/configs/emulation/full/permutation/rank6_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/emulation/full/permutation/rank6_test_params.yaml @@ -1,9 +1,12 @@ --- Log Level: [ HIPTENSOR_LOG_LEVEL_ERROR, HIPTENSOR_LOG_LEVEL_PERF_TRACE ] Tensor Data Types: + - [ HIPTENSOR_R_64F, HIPTENSOR_R_64F] - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_32F] Alphas: - 2.3 Lengths: diff --git a/projects/hiptensor/test/02_elementwise/configs/emulation/quick/permutation/rank2_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/emulation/quick/permutation/rank2_test_params.yaml index cfeebd682618..3de552fa0308 100644 --- a/projects/hiptensor/test/02_elementwise/configs/emulation/quick/permutation/rank2_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/emulation/quick/permutation/rank2_test_params.yaml @@ -4,6 +4,7 @@ Tensor Data Types: - [ HIPTENSOR_R_64F, HIPTENSOR_R_64F] - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] Alphas: - 2.3 Lengths: diff --git a/projects/hiptensor/test/02_elementwise/configs/emulation/quick/permutation/rank3_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/emulation/quick/permutation/rank3_test_params.yaml index 429787b9591e..bd2da8fd054c 100644 --- a/projects/hiptensor/test/02_elementwise/configs/emulation/quick/permutation/rank3_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/emulation/quick/permutation/rank3_test_params.yaml @@ -4,6 +4,7 @@ Tensor Data Types: - [ HIPTENSOR_R_64F, HIPTENSOR_R_64F] - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] Alphas: - 2.3 Lengths: diff --git a/projects/hiptensor/test/02_elementwise/configs/emulation/quick/permutation/rank4_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/emulation/quick/permutation/rank4_test_params.yaml index 4396d8d2443c..c5115f1fa33a 100644 --- a/projects/hiptensor/test/02_elementwise/configs/emulation/quick/permutation/rank4_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/emulation/quick/permutation/rank4_test_params.yaml @@ -4,6 +4,7 @@ Tensor Data Types: - [ HIPTENSOR_R_64F, HIPTENSOR_R_64F] - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] Alphas: - 2.3 Lengths: diff --git a/projects/hiptensor/test/02_elementwise/configs/emulation/quick/permutation/rank5_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/emulation/quick/permutation/rank5_test_params.yaml index 8b1b2e46ade4..db5feee90973 100644 --- a/projects/hiptensor/test/02_elementwise/configs/emulation/quick/permutation/rank5_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/emulation/quick/permutation/rank5_test_params.yaml @@ -4,6 +4,7 @@ Tensor Data Types: - [ HIPTENSOR_R_64F, HIPTENSOR_R_64F] - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] Alphas: - 2.3 Lengths: diff --git a/projects/hiptensor/test/02_elementwise/configs/emulation/quick/permutation/rank6_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/emulation/quick/permutation/rank6_test_params.yaml index 74414f43fdbf..dd817f4c33fd 100644 --- a/projects/hiptensor/test/02_elementwise/configs/emulation/quick/permutation/rank6_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/emulation/quick/permutation/rank6_test_params.yaml @@ -4,6 +4,7 @@ Tensor Data Types: - [ HIPTENSOR_R_64F, HIPTENSOR_R_64F] - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] Alphas: - 2.3 Lengths: diff --git a/projects/hiptensor/test/02_elementwise/configs/validation/comprehensive/permutation/rank2_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/validation/comprehensive/permutation/rank2_test_params.yaml index 0bd39adc4471..0afb7fbcc5c1 100644 --- a/projects/hiptensor/test/02_elementwise/configs/validation/comprehensive/permutation/rank2_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/validation/comprehensive/permutation/rank2_test_params.yaml @@ -5,6 +5,8 @@ Tensor Data Types: - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_32F] Alphas: - 0.0 - 1.0 diff --git a/projects/hiptensor/test/02_elementwise/configs/validation/comprehensive/permutation/rank3_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/validation/comprehensive/permutation/rank3_test_params.yaml index 86eab99b7a44..19996e9fd8b5 100644 --- a/projects/hiptensor/test/02_elementwise/configs/validation/comprehensive/permutation/rank3_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/validation/comprehensive/permutation/rank3_test_params.yaml @@ -5,6 +5,8 @@ Tensor Data Types: - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_32F] Alphas: - 0.0 - 1.0 diff --git a/projects/hiptensor/test/02_elementwise/configs/validation/comprehensive/permutation/rank4_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/validation/comprehensive/permutation/rank4_test_params.yaml index b482ef0ccd22..1d404d62cf44 100644 --- a/projects/hiptensor/test/02_elementwise/configs/validation/comprehensive/permutation/rank4_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/validation/comprehensive/permutation/rank4_test_params.yaml @@ -5,6 +5,8 @@ Tensor Data Types: - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_32F] Alphas: - 0.0 - 1.0 diff --git a/projects/hiptensor/test/02_elementwise/configs/validation/comprehensive/permutation/rank5_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/validation/comprehensive/permutation/rank5_test_params.yaml index 6eb2c607556b..171608cfa79a 100644 --- a/projects/hiptensor/test/02_elementwise/configs/validation/comprehensive/permutation/rank5_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/validation/comprehensive/permutation/rank5_test_params.yaml @@ -5,6 +5,8 @@ Tensor Data Types: - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_32F] Alphas: - 0.0 - 1.0 diff --git a/projects/hiptensor/test/02_elementwise/configs/validation/comprehensive/permutation/rank6_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/validation/comprehensive/permutation/rank6_test_params.yaml index 126cedd0f806..22b13b39b5c8 100644 --- a/projects/hiptensor/test/02_elementwise/configs/validation/comprehensive/permutation/rank6_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/validation/comprehensive/permutation/rank6_test_params.yaml @@ -5,6 +5,8 @@ Tensor Data Types: - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_32F] Alphas: - 0.0 - 2.3 diff --git a/projects/hiptensor/test/02_elementwise/configs/validation/quick/permutation/rank2_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/validation/quick/permutation/rank2_test_params.yaml index cfeebd682618..3de552fa0308 100644 --- a/projects/hiptensor/test/02_elementwise/configs/validation/quick/permutation/rank2_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/validation/quick/permutation/rank2_test_params.yaml @@ -4,6 +4,7 @@ Tensor Data Types: - [ HIPTENSOR_R_64F, HIPTENSOR_R_64F] - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] Alphas: - 2.3 Lengths: diff --git a/projects/hiptensor/test/02_elementwise/configs/validation/quick/permutation/rank3_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/validation/quick/permutation/rank3_test_params.yaml index 429787b9591e..bd2da8fd054c 100644 --- a/projects/hiptensor/test/02_elementwise/configs/validation/quick/permutation/rank3_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/validation/quick/permutation/rank3_test_params.yaml @@ -4,6 +4,7 @@ Tensor Data Types: - [ HIPTENSOR_R_64F, HIPTENSOR_R_64F] - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] Alphas: - 2.3 Lengths: diff --git a/projects/hiptensor/test/02_elementwise/configs/validation/quick/permutation/rank4_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/validation/quick/permutation/rank4_test_params.yaml index 4396d8d2443c..c5115f1fa33a 100644 --- a/projects/hiptensor/test/02_elementwise/configs/validation/quick/permutation/rank4_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/validation/quick/permutation/rank4_test_params.yaml @@ -4,6 +4,7 @@ Tensor Data Types: - [ HIPTENSOR_R_64F, HIPTENSOR_R_64F] - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] Alphas: - 2.3 Lengths: diff --git a/projects/hiptensor/test/02_elementwise/configs/validation/quick/permutation/rank5_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/validation/quick/permutation/rank5_test_params.yaml index 8b1b2e46ade4..db5feee90973 100644 --- a/projects/hiptensor/test/02_elementwise/configs/validation/quick/permutation/rank5_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/validation/quick/permutation/rank5_test_params.yaml @@ -4,6 +4,7 @@ Tensor Data Types: - [ HIPTENSOR_R_64F, HIPTENSOR_R_64F] - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] Alphas: - 2.3 Lengths: diff --git a/projects/hiptensor/test/02_elementwise/configs/validation/quick/permutation/rank6_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/validation/quick/permutation/rank6_test_params.yaml index 74414f43fdbf..dd817f4c33fd 100644 --- a/projects/hiptensor/test/02_elementwise/configs/validation/quick/permutation/rank6_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/validation/quick/permutation/rank6_test_params.yaml @@ -4,6 +4,7 @@ Tensor Data Types: - [ HIPTENSOR_R_64F, HIPTENSOR_R_64F] - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] Alphas: - 2.3 Lengths: diff --git a/projects/hiptensor/test/02_elementwise/configs/validation/standard/permutation/rank2_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/validation/standard/permutation/rank2_test_params.yaml index 3cdd8dd32831..b38594729651 100644 --- a/projects/hiptensor/test/02_elementwise/configs/validation/standard/permutation/rank2_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/validation/standard/permutation/rank2_test_params.yaml @@ -5,6 +5,8 @@ Tensor Data Types: - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_32F] Alphas: - 0.0 - 1.0 diff --git a/projects/hiptensor/test/02_elementwise/configs/validation/standard/permutation/rank3_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/validation/standard/permutation/rank3_test_params.yaml index 8b85fcd0b00e..221e76a81b6e 100644 --- a/projects/hiptensor/test/02_elementwise/configs/validation/standard/permutation/rank3_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/validation/standard/permutation/rank3_test_params.yaml @@ -5,6 +5,8 @@ Tensor Data Types: - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_32F] Alphas: - 0.0 - 1.0 diff --git a/projects/hiptensor/test/02_elementwise/configs/validation/standard/permutation/rank4_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/validation/standard/permutation/rank4_test_params.yaml index ce45b14d4a1d..f2ae34579f8b 100644 --- a/projects/hiptensor/test/02_elementwise/configs/validation/standard/permutation/rank4_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/validation/standard/permutation/rank4_test_params.yaml @@ -5,6 +5,8 @@ Tensor Data Types: - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_32F] Alphas: - 0.0 - 1.0 diff --git a/projects/hiptensor/test/02_elementwise/configs/validation/standard/permutation/rank5_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/validation/standard/permutation/rank5_test_params.yaml index 31d650c697cc..075f6ae73367 100644 --- a/projects/hiptensor/test/02_elementwise/configs/validation/standard/permutation/rank5_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/validation/standard/permutation/rank5_test_params.yaml @@ -5,6 +5,8 @@ Tensor Data Types: - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_32F] Alphas: - 0.0 - 1.0 diff --git a/projects/hiptensor/test/02_elementwise/configs/validation/standard/permutation/rank6_test_params.yaml b/projects/hiptensor/test/02_elementwise/configs/validation/standard/permutation/rank6_test_params.yaml index 4351c638ce7f..39a1f590ff08 100644 --- a/projects/hiptensor/test/02_elementwise/configs/validation/standard/permutation/rank6_test_params.yaml +++ b/projects/hiptensor/test/02_elementwise/configs/validation/standard/permutation/rank6_test_params.yaml @@ -5,6 +5,8 @@ Tensor Data Types: - [ HIPTENSOR_R_32F, HIPTENSOR_R_32F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_16F] - [ HIPTENSOR_R_16F, HIPTENSOR_R_32F] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_16BF] + - [ HIPTENSOR_R_16BF, HIPTENSOR_R_32F] Alphas: - 0.0 - 2.3 diff --git a/projects/hiptensor/test/02_elementwise/elementwise_cpu_impl_test.cpp b/projects/hiptensor/test/02_elementwise/elementwise_cpu_impl_test.cpp index c784cdcf12b3..5251227cfb20 100644 --- a/projects/hiptensor/test/02_elementwise/elementwise_cpu_impl_test.cpp +++ b/projects/hiptensor/test/02_elementwise/elementwise_cpu_impl_test.cpp @@ -35,7 +35,8 @@ template auto permuteWithCpu(hiptensorDataType_t typeA, hiptensorDataType_t typeB, - hiptensorDataType_t typeCompute) + hiptensorDataType_t typeCompute, + double tolerance = 0.0) { std::vector modeA{'w', 'h', 'c', 'n'}; std::vector modeB{'c', 'n', 'h', 'w'}; @@ -79,15 +80,15 @@ auto permuteWithCpu(hiptensorDataType_t typeA, std::vector aArray(elementsA); std::vector bArray(elementsB); - std::iota(aArray.begin(), aArray.end(), 0); - std::vector referenceArray; + std::iota(aArray.begin(), aArray.end(), static_cast(0.0f)); + std::vector referenceValues; using hiptensor::HiptensorOptions; auto& options = HiptensorOptions::instance(); if(options->isColMajorStrides()) { - referenceArray = { + referenceValues = { 0., -12.6, -25.2, -37.8, -50.4, -63., -75.6, -88.2, -100.8, -113.4, -126., -138.6, -151.2, -163.8, -176.4, -189., -201.6, -214.2, -226.8, -239.4, -6.3, -18.9, -31.5, -44.1, -56.7, -69.3, -81.9, -94.5, -107.1, -119.7, -132.3, -144.9, -157.5, @@ -102,7 +103,7 @@ auto permuteWithCpu(hiptensorDataType_t typeA, } else { - referenceArray = { + referenceValues = { 0., -84., -168., -42., -126., -210., -2.1, -86.1, -170.1, -44.1, -128.1, -212.1, -4.2, -88.2, -172.2, -46.2, -130.2, -214.2, -6.3, -90.3, -174.3, -48.3, -132.3, -216.3, -8.4, -92.4, -176.4, -50.4, -134.4, -218.4, -10.5, -94.5, -178.5, @@ -116,7 +117,13 @@ auto permuteWithCpu(hiptensorDataType_t typeA, -205.8, -79.8, -163.8, -247.8, -39.9, -123.9, -207.9, -81.9, -165.9, -249.9}; } - const floatTypeCompute alphaValue = 2.1f; + std::vector referenceArray; + for(auto v : referenceValues) + { + referenceArray.push_back(static_cast(v)); + } + + const floatTypeCompute alphaValue = static_cast(2.1); hiptensorHandle_t handle; CHECK_HIPTENSOR_ERROR(hiptensorCreate(&handle)); hiptensorTensorDescriptor_t descA = nullptr; @@ -153,7 +160,7 @@ auto permuteWithCpu(hiptensorDataType_t typeA, bArray.data(), bArray.size(), hiptensor::convertToComputeType(typeCompute), - 0); + tolerance); } TEST(PermutationCpuImplTest, CompareF32ResultWithReference) @@ -185,3 +192,54 @@ TEST(PermutationCpuImplTest, CompareF16ResultWithReference) = permuteWithCpu(typeA, typeB, typeCompute); EXPECT_TRUE(result) << "max_relative_error: " << maxRelativeError; } + +TEST(PermutationCpuImplTest, CompareF64ResultWithReference) +{ + typedef double floatTypeA; + typedef double floatTypeB; + typedef double floatTypeCompute; + + hiptensorDataType_t typeA = HIPTENSOR_R_64F; + hiptensorDataType_t typeB = HIPTENSOR_R_64F; + hiptensorDataType_t typeCompute = HIPTENSOR_R_64F; + + if(!isF64Supported()) + { + GTEST_SKIP(); + } + + // Tight enough that a reference computing in float (~1e-8 relative error) fails. + auto [result, maxRelativeError] = permuteWithCpu( + typeA, typeB, typeCompute, 1e-12); + EXPECT_TRUE(result) << "max_relative_error: " << maxRelativeError; +} + +TEST(PermutationCpuImplTest, CompareBF16ResultWithReference) +{ + typedef hip_bfloat16 floatTypeA; + typedef hip_bfloat16 floatTypeB; + typedef hip_bfloat16 floatTypeCompute; + + hiptensorDataType_t typeA = HIPTENSOR_R_16BF; + hiptensorDataType_t typeB = HIPTENSOR_R_16BF; + hiptensorDataType_t typeCompute = HIPTENSOR_R_16BF; + + auto [result, maxRelativeError] + = permuteWithCpu(typeA, typeB, typeCompute); + EXPECT_TRUE(result) << "max_relative_error: " << maxRelativeError; +} + +TEST(PermutationCpuImplTest, CompareBF16WithF32ScalarResultWithReference) +{ + typedef hip_bfloat16 floatTypeA; + typedef hip_bfloat16 floatTypeB; + typedef float floatTypeCompute; + + hiptensorDataType_t typeA = HIPTENSOR_R_16BF; + hiptensorDataType_t typeB = HIPTENSOR_R_16BF; + hiptensorDataType_t typeCompute = HIPTENSOR_R_32F; + + auto [result, maxRelativeError] + = permuteWithCpu(typeA, typeB, typeCompute); + EXPECT_TRUE(result) << "max_relative_error: " << maxRelativeError; +} diff --git a/projects/hiptensor/test/02_elementwise/elementwise_resource.cpp b/projects/hiptensor/test/02_elementwise/elementwise_resource.cpp index 662ea1b44550..1996ef7140ec 100644 --- a/projects/hiptensor/test/02_elementwise/elementwise_resource.cpp +++ b/projects/hiptensor/test/02_elementwise/elementwise_resource.cpp @@ -149,6 +149,11 @@ namespace hiptensor { fillLaunchKernel((float*)devicePtr.get(), mCurrentMatrixElement, seed); } + else if(mCurrentDataType == HIPTENSOR_R_16BF) + { + fillLaunchKernel( + (hip_bfloat16*)devicePtr.get(), mCurrentMatrixElement, seed); + } else { fillLaunchKernel<_Float16>((_Float16*)devicePtr.get(), mCurrentMatrixElement, seed); diff --git a/projects/hiptensor/test/02_elementwise/elementwise_test.cpp b/projects/hiptensor/test/02_elementwise/elementwise_test.cpp index 901adb7574c1..e72fcaf16c7e 100644 --- a/projects/hiptensor/test/02_elementwise/elementwise_test.cpp +++ b/projects/hiptensor/test/02_elementwise/elementwise_test.cpp @@ -58,7 +58,7 @@ namespace hiptensor // False = skip test bool PermutationTest::checkDevice(hiptensorDataType_t datatype) const { - return (isF16Supported() && (datatype == HIPTENSOR_R_16F)) + return (isF16Supported() && (datatype == HIPTENSOR_R_16F || datatype == HIPTENSOR_R_16BF)) || (isF32Supported() && (datatype == HIPTENSOR_R_32F)) || (isF64Supported() && (datatype == HIPTENSOR_R_64F)); } @@ -178,10 +178,10 @@ namespace hiptensor auto op = operators[0]; EXPECT_TRUE((op == HIPTENSOR_OP_IDENTITY) || (op == HIPTENSOR_OP_NEG)); - EXPECT_EQ(dataTypes.size(), 2); // HIPTENSOR_R_16F, HIPTENSOR_R_32F or HIPTENSOR_R_64F + EXPECT_EQ(dataTypes.size(), 2); // {data type, compute type} auto abDataType = dataTypes[0]; - EXPECT_TRUE((abDataType == HIPTENSOR_R_16F) || (abDataType == HIPTENSOR_R_32F) - || (abDataType == HIPTENSOR_R_64F)); + EXPECT_TRUE((abDataType == HIPTENSOR_R_16F) || (abDataType == HIPTENSOR_R_16BF) + || (abDataType == HIPTENSOR_R_32F) || (abDataType == HIPTENSOR_R_64F)); mRunFlag &= checkDevice(abDataType); @@ -242,6 +242,40 @@ namespace hiptensor stream, (float*)resource->hostReference().get(), elementsRef); stream << std::endl; } + else if(dataType == HIPTENSOR_R_16BF) + { + stream << "Tensor A elements (" << elementsA << "):\n"; + hiptensorPrintArrayElements( + stream, (hip_bfloat16*)resource->hostInput1().get(), elementsA); + stream << std::endl; + + stream << "Tensor B elements (" << elementsB << "):\n"; + hiptensorPrintArrayElements( + stream, (hip_bfloat16*)resource->hostOutput().get(), elementsB); + stream << std::endl; + + stream << "Tensor ref elements (" << elementsRef << "):\n"; + hiptensorPrintArrayElements( + stream, (hip_bfloat16*)resource->hostReference().get(), elementsRef); + stream << std::endl; + } + else if(dataType == HIPTENSOR_R_64F) + { + stream << "Tensor A elements (" << elementsA << "):\n"; + hiptensorPrintArrayElements( + stream, (double*)resource->hostInput1().get(), elementsA); + stream << std::endl; + + stream << "Tensor B elements (" << elementsB << "):\n"; + hiptensorPrintArrayElements( + stream, (double*)resource->hostOutput().get(), elementsB); + stream << std::endl; + + stream << "Tensor ref elements (" << elementsRef << "):\n"; + hiptensorPrintArrayElements( + stream, (double*)resource->hostReference().get(), elementsRef); + stream << std::endl; + } else { stream << "Tensor A elements (" << elementsA << "):\n"; @@ -389,6 +423,11 @@ namespace hiptensor { *(reinterpret_cast<_Float16*>(&alphaValue)) = static_cast<_Float16>(alpha); } + else if(computeDataType == HIPTENSOR_R_16BF) + { + *(reinterpret_cast(&alphaValue)) + = static_cast(static_cast(alpha)); + } else if(computeDataType == HIPTENSOR_R_32F) { *(reinterpret_cast(&alphaValue)) = static_cast(alpha); @@ -502,6 +541,29 @@ namespace hiptensor resource->getCurrentMatrixElement(), convertToComputeType(computeDataType)); } + else if(abDataType == HIPTENSOR_R_16BF) + { + CHECK_HIPTENSOR_ERROR(hiptensorElementwisePermuteReference( + &alphaValue, + (const hip_bfloat16*)resource->hostInput1().get(), + descA, + modeA.data(), + Aop, + (hip_bfloat16*)resource->hostReference().get(), + descB, + modeB.data(), + computeDataType, + 0 /* stream */)); + + resource->copyReferenceToDevice(); + + std::tie(mValidationResult, mMaxRelativeError) + = compareEqualLaunchKernel( + (hip_bfloat16*)resource->deviceOutput().get(), + (hip_bfloat16*)resource->deviceReference().get(), + resource->getCurrentMatrixElement(), + convertToComputeType(computeDataType)); + } else if(abDataType == HIPTENSOR_R_64F) { CHECK_HIPTENSOR_ERROR(hiptensorElementwisePermuteReference( diff --git a/projects/hiptensor/test/utils.hpp b/projects/hiptensor/test/utils.hpp index ecf7eef8af84..e46815022cb9 100644 --- a/projects/hiptensor/test/utils.hpp +++ b/projects/hiptensor/test/utils.hpp @@ -375,8 +375,16 @@ std::pair compareEqual(DDataType const* deviceD, bool retval = true; double max_relative_error = 0.0; - auto toDouble - = [](DDataType const& val) { return static_cast(static_cast(val)); }; + auto toDouble = [](DDataType const& val) { + if constexpr(std::is_same_v) + { + return val; + } + else + { + return static_cast(static_cast(val)); + } + }; bool isInf = false; bool isNaN = false;