diff --git a/.github/copilot-instructions.md b/.github/copilot-instructions.md index 57fdb035..8134f9b6 100644 --- a/.github/copilot-instructions.md +++ b/.github/copilot-instructions.md @@ -79,7 +79,7 @@ cmake --build build --config Release --parallel 32 - **CUDA >= 13**: All architectures allowed. ### Windows-Specific Notes -- CI uses self-hosted runners with LLVM 21.1.0 at `D:/clang+llvm-21.1.0-x86_64-pc-windows-msvc/bin/`. +- CI uses self-hosted runners with LLVM 23.1.0 at `D:/clang+llvm-23.1.0-x86_64-pc-windows-msvc/bin/`. - The Ninja generator requires a workaround: after `cmake`, `rules.ninja` may contain a wrong NVCC path that must be patched (see `cmake-windows-amd64.yml`). - The VS Developer Shell (`Enter-VsDevShell`) must be activated for both configure and build steps on Windows. - `utf8cp.manifest` is embedded into test executables on Windows for UTF-8 codepage support. @@ -222,9 +222,9 @@ Each operation type exposes a `static build(...)` factory method returning an in All three workflow files trigger on **pull requests to `main`** (push triggers are commented out). All runners are **self-hosted**. ### Linux (cmake-linux-amd64.yml, cmake-linux-arm64.yml) -- **Compilers**: `g++-13`, `clang++-21` +- **Compilers**: `g++-13`, `clang++-23` - **CUDA**: 13.3 (via `/usr/local/cuda-/bin/nvcc`) -- **CMake**: Custom installation at `/home/cudeiro/cmake-4.3.3-linux-x86_64/bin/` (added to PATH) +- **CMake**: Custom installation at `/home/cudeiro/cmake-4.4.2-linux-x86_64/bin/` (added to PATH) - **Generator**: Ninja - **Build type**: Release @@ -232,7 +232,7 @@ All three workflow files trigger on **pull requests to `main`** (push triggers a - **Host compilers**: `cl` (MSVC), `clang-cl` - **MSVC versions**: 14.44, 14.51 (via `-vcvars_ver`) - **CUDA**: 13.0, 13.3 (NVCC at `%ProgramFiles%\NVIDIA GPU Computing Toolkit\CUDA\v\bin\nvcc.exe`) -- **LLVM**: `D:/clang+llvm-21.1.0-x86_64-pc-windows-msvc/bin/` (added to PATH) +- **LLVM**: `D:/clang+llvm-23.1.0-x86_64-pc-windows-msvc/bin/` (added to PATH) - **Generator**: Ninja - **Workaround**: After CMake configure, `rules.ninja` may contain an empty NVCC path that is patched with PowerShell string replacement. diff --git a/.github/skills/fkl-build-and-test/SKILL.md b/.github/skills/fkl-build-and-test/SKILL.md index 0685ced5..559edb51 100644 --- a/.github/skills/fkl-build-and-test/SKILL.md +++ b/.github/skills/fkl-build-and-test/SKILL.md @@ -102,7 +102,7 @@ int launch() { RUN_ALL_TESTS } ## CI Matrix GitHub workflows build on `linux-amd64`, `linux-arm64` and `windows-amd64` using self-hosted runners. -- Linux builds against `g++-13` and `clang++-21`. +- Linux builds against `g++-13` and `clang++-23`. - Windows builds against MSVC (`cl` versions 14.44 and 14.51) and `clang-cl`. - Keep changes warning-clean on BOTH compilers: clang is a first-class compiler for FKL (single-step host+device compiles matter for downstream packaging). diff --git a/.github/workflows/cmake-linux-arm64.yml b/.github/workflows/cmake-linux-arm64.yml index 6f52ec19..d2c46fa2 100644 --- a/.github/workflows/cmake-linux-arm64.yml +++ b/.github/workflows/cmake-linux-arm64.yml @@ -1,4 +1,4 @@ -name: CMake on Linux (ARM64) +name: CMake on Linux (CUDA ARM64) on: #push: @@ -19,7 +19,7 @@ jobs: include: - host_compiler: "g++-13" cuda_toolkit: "13.3" - - host_compiler: "clang++-21" + - host_compiler: "clang++-23" cuda_toolkit: "13.3" @@ -30,14 +30,14 @@ jobs: id: strings run: | echo "build-output-dir=${{github.workspace}}/build" >> "$GITHUB_OUTPUT" - echo "PATH=$HOME/cmake-4.3.3-linux-aarch64/bin/:$PATH" >> "$GITHUB_ENV" + echo "PATH=$HOME/cmake-4.4.2-linux-aarch64/bin/:$PATH" >> "$GITHUB_ENV" echo "CUDACXX=/usr/local/cuda-${{matrix.cuda_toolkit}}/bin/nvcc" >> "$GITHUB_ENV" echo "CC=${{matrix.host_compiler}}" >> "$GITHUB_ENV" echo "CXX=${{matrix.host_compiler}}" >> "$GITHUB_ENV" - name: Configure CMake run: | - cmake -G "Ninja" -B ${{steps.strings.outputs.build-output-dir}} -DCMAKE_BUILD_TYPE="Release" -S ${{github.workspace}} + cmake -G "Ninja" -DENABLE_HIP=OFF -DENABLE_CUDA=ON -B ${{steps.strings.outputs.build-output-dir}} -DCMAKE_BUILD_TYPE="Release" -S ${{github.workspace}} - name: Build run: | diff --git a/.github/workflows/cmake-linux-amd64.yml b/.github/workflows/cmake-linux-cuda-amd64.yml similarity index 82% rename from .github/workflows/cmake-linux-amd64.yml rename to .github/workflows/cmake-linux-cuda-amd64.yml index ff90f5af..ebcc24f5 100644 --- a/.github/workflows/cmake-linux-amd64.yml +++ b/.github/workflows/cmake-linux-cuda-amd64.yml @@ -1,4 +1,4 @@ -name: CMake on Linux (x86_64) +name: CMake on Linux (CUDA x86_64) on: #push: @@ -19,7 +19,7 @@ jobs: include: - host_compiler: "g++-13" cuda_toolkit: "13.3" - - host_compiler: "clang++-21" + - host_compiler: "clang++-23" cuda_toolkit: "13.3" @@ -30,13 +30,13 @@ jobs: id: strings run: | echo "build-output-dir=${{github.workspace}}/build" >> "$GITHUB_OUTPUT" - echo "PATH=$HOME/cmake-4.3.3-linux-x86_64/bin/:$PATH" >> "$GITHUB_ENV" + echo "PATH=$HOME/cmake-4.4.2-linux-x86_64/bin/:$PATH" >> "$GITHUB_ENV" echo "CUDACXX=/usr/local/cuda-${{matrix.cuda_toolkit}}/bin/nvcc" >> "$GITHUB_ENV" echo "CC=${{matrix.host_compiler}}" >> "$GITHUB_ENV" echo "CXX=${{matrix.host_compiler}}" >> "$GITHUB_ENV" - name: Configure CMake run: | - cmake -G "Ninja" -B ${{steps.strings.outputs.build-output-dir}} -DCMAKE_BUILD_TYPE="Release" -S ${{github.workspace}} + cmake -G "Ninja" -DENABLE_HIP=OFF -DENABLE_CUDA=ON -B ${{steps.strings.outputs.build-output-dir}} -DCMAKE_BUILD_TYPE="Release" -S ${{github.workspace}} - name: Build run: | diff --git a/.github/workflows/cmake-linux-rocm-amd64.yml b/.github/workflows/cmake-linux-rocm-amd64.yml new file mode 100644 index 00000000..aef379bf --- /dev/null +++ b/.github/workflows/cmake-linux-rocm-amd64.yml @@ -0,0 +1,49 @@ +name: CMake on Linux (ROCM x86_64) + +on: + #push: + # branches: [ $default-branch ] + pull_request: + branches: ["main"] + +jobs: + build: + runs-on: [self-hosted, linux, x64, rocm] + + strategy: + # Set fail-fast to false to ensure that feedback is delivered for all matrix combinations. + # Consider changing this to true when your workflow is stable. + fail-fast: false + + matrix: + include: + - rocm_version: "10.0" + + steps: + - uses: actions/checkout@v6 + - name: Set reusable strings and path + # Turn repeated input strings (such as the build output directory) into step outputs. These step outputs can be used throughout the workflow file. + id: strings + run: | + echo "build-output-dir=${{github.workspace}}/build" >> "$GITHUB_OUTPUT" + ROCM_PATH=/opt/rocm/core-${{matrix.rocm_version}} + echo "ROCM_PATH=$ROCM_PATH" >> "$GITHUB_ENV" + echo "HIP_PATH=$ROCM_PATH" >> "$GITHUB_ENV" + echo "PATH=$HOME/cmake-4.4.2-linux-x86_64/bin/:$ROCM_PATH/bin:$PATH" >> "$GITHUB_ENV" + echo "CC=/opt/rocm/bin/amdclang" >> "$GITHUB_ENV" + echo "CXX=/opt/rocm/bin/amdclang++" >> "$GITHUB_ENV" + echo "HIPCXX=/opt/rocm/bin/amdclang++" >> "$GITHUB_ENV" + echo "LD_LIBRARY_PATH=/opt/rocm/core-${{matrix.rocm_version}}/lib:${LD_LIBRARY_PATH}" >> "$GITHUB_ENV" + + - name: Configure CMake + run: | + cmake -G "Ninja" -DENABLE_HIP=ON -DENABLE_CUDA=OFF -B ${{steps.strings.outputs.build-output-dir}} -DCMAKE_BUILD_TYPE="Release" -S ${{github.workspace}} + + - name: Build + run: | + cmake --build ${{ steps.strings.outputs.build-output-dir }} --config Release + + - name: Test + working-directory: ${{ steps.strings.outputs.build-output-dir }} + run: | + ctest --build-config Release --output-junit test_results.xml diff --git a/.github/workflows/cmake-windows-amd64.yml b/.github/workflows/cmake-windows-cuda-amd64.yml similarity index 92% rename from .github/workflows/cmake-windows-amd64.yml rename to .github/workflows/cmake-windows-cuda-amd64.yml index 5b62ccd0..6897fc76 100644 --- a/.github/workflows/cmake-windows-amd64.yml +++ b/.github/workflows/cmake-windows-cuda-amd64.yml @@ -1,4 +1,4 @@ -name: CMake on Windows (x64) +name: CMake on Windows (CUDA x64) on: #push: @@ -35,7 +35,7 @@ jobs: id: strings run: | echo "build-output-dir=${{github.workspace}}/build" >> "$env:GITHUB_OUTPUT" - echo "D:/clang+llvm-21.1.0-x86_64-pc-windows-msvc/bin/" | Out-File -FilePath $env:GITHUB_PATH -Encoding utf8 -Append + echo "D:/clang+llvm-23.1.0-x86_64-pc-windows-msvc/bin/" | Out-File -FilePath $env:GITHUB_PATH -Encoding utf8 -Append echo "CUDACXX=${env:ProgramFiles}\NVIDIA GPU Computing Toolkit\CUDA\v${{matrix.cuda_version}}\bin\nvcc.exe" >> "$env:GITHUB_ENV" echo "CC=${{matrix.host_compiler}}" >> "$env:GITHUB_ENV" echo "CXX=${{matrix.host_compiler}}" >> "$env:GITHUB_ENV" @@ -51,7 +51,7 @@ jobs: run: | Import-Module (Get-ChildItem $env:VS_INSTALLATION_PATH -Recurse -File -Filter Microsoft.VisualStudio.DevShell.dll).FullName Enter-VsDevShell -VsInstallPath $env:VS_INSTALLATION_PATH -SkipAutomaticLocation -DevCmdArguments $env:DEV_ARGS - cmake -G "Ninja" -B ${{steps.strings.outputs.build-output-dir}} -DCMAKE_BUILD_TYPE=Release -S ${{github.workspace}}; + cmake -G "Ninja" -DENABLE_HIP=OFF -DENABLE_CUDA=ON -B ${{steps.strings.outputs.build-output-dir}} -DCMAKE_BUILD_TYPE=Release -S ${{github.workspace}}; if ($LASTEXITCODE -ne 0) { exit $LASTEXITCODE } #workaround for rules.ninja generating empty path (Get-Content ${{steps.strings.outputs.build-output-dir}}\CMakeFiles\rules.ninja) -replace "\\nvcc\\bin\\nvcc.exe", "$env:CUDACXX" | Set-Content ${{steps.strings.outputs.build-output-dir}}\CMakeFiles\rules.ninja diff --git a/.github/workflows/cmake-windows-rocm-amd64.yml.disabled b/.github/workflows/cmake-windows-rocm-amd64.yml.disabled new file mode 100644 index 00000000..c1d5cb3f --- /dev/null +++ b/.github/workflows/cmake-windows-rocm-amd64.yml.disabled @@ -0,0 +1,59 @@ +name: CMake on Windows (ROCM x64) + +on: + #push: + # branches: [ $default-branch ] + pull_request: + branches: ["main"] + +jobs: + build: + runs-on: [self-hosted, windows, x64, rocm] + + strategy: + # Set fail-fast to false to ensure that feedback is delivered for all matrix combinations. Consider changing this to true when your workflow is stable. + fail-fast: false + matrix: + include: + - msvc_env_version: "14.44" #HIP on windows is vs2022 only + rocm_version: "7.2" + + + steps: + - uses: actions/checkout@v6 + + - name: Set reusable strings + # Turn repeated input strings (such as the build output directory) into step outputs. These step outputs can be used throughout the workflow file. + id: strings + run: | + echo "build-output-dir=${{github.workspace}}/build" >> "$env:GITHUB_OUTPUT" + echo "CC=$env:HIP_PATH\bin\clang.exe" >> "$env:GITHUB_ENV" + echo "CXX=$env:HIP_PATH\bin\clang++.exe" >> "$env:GITHUB_ENV" + echo "HIPCXX=$env:HIP_PATH\bin\clang++.exe" >> "$env:GITHUB_ENV" + echo "CMAKE_PREFIX_PATH=$env:HIP_PATH" >> "$env:GITHUB_ENV" + + - name: Detect visual studio installation (2022) + run: | + $vsWhere = "${Env:ProgramFiles(x86)}/Microsoft Visual Studio/Installer/vswhere.exe" + $vsInstallationPath = & $vsWhere -products * -version "[17.0,18.0)" -property installationPath + echo "VS_INSTALLATION_PATH=$vsInstallationPath" >> "$env:GITHUB_ENV" + echo "DEV_ARGS=-arch=x64 -vcvars_ver=${{matrix.msvc_env_version}}" >> "$env:GITHUB_ENV" + + - name: Configure CMake + run: | + Import-Module (Get-ChildItem $env:VS_INSTALLATION_PATH -Recurse -File -Filter Microsoft.VisualStudio.DevShell.dll).FullName + Enter-VsDevShell -VsInstallPath $env:VS_INSTALLATION_PATH -SkipAutomaticLocation -DevCmdArguments $env:DEV_ARGS + cmake -G "Ninja" -DENABLE_HIP=ON -DENABLE_CUDA=OFF ` + -DCMAKE_HIP_ARCHITECTURES="gfx1200" -B ${{steps.strings.outputs.build-output-dir}} -DCMAKE_BUILD_TYPE=Release -S ${{github.workspace}}; + if ($LASTEXITCODE -ne 0) { exit $LASTEXITCODE } + + + - name: Build + run: | + Import-Module (Get-ChildItem $env:VS_INSTALLATION_PATH -Recurse -File -Filter Microsoft.VisualStudio.DevShell.dll).FullName + Enter-VsDevShell -VsInstallPath $env:VS_INSTALLATION_PATH -SkipAutomaticLocation -DevCmdArguments $env:DEV_ARGS + cmake --build ${{ steps.strings.outputs.build-output-dir }} --config Release + + - name: Test + working-directory: ${{ steps.strings.outputs.build-output-dir }} + run: ctest --build-config Release --output-junit test_results.xml diff --git a/.github/workflows/copilot-setup-steps.yml b/.github/workflows/copilot-setup-steps.yml index 0b8e1084..dfe41e67 100644 --- a/.github/workflows/copilot-setup-steps.yml +++ b/.github/workflows/copilot-setup-steps.yml @@ -17,7 +17,7 @@ jobs: - name: Set up Copilot Environment # The agent will need the same paths you define in your CI run: | - echo "PATH=$HOME/cmake-4.3.3-linux-x86_64/bin/:$PATH" >> "$GITHUB_ENV" + echo "PATH=$HOME/cmake-4.4.2-linux-x86_64/bin/:$PATH" >> "$GITHUB_ENV" echo "CUDACXX=/usr/local/cuda-13.3/bin/nvcc" >> "$GITHUB_ENV" echo "CC=g++-13" >> "$GITHUB_ENV" echo "CXX=g++-13" >> "$GITHUB_ENV" \ No newline at end of file diff --git a/CLAUDE.md b/CLAUDE.md index 843df7cc..fc154841 100644 --- a/CLAUDE.md +++ b/CLAUDE.md @@ -29,7 +29,7 @@ ctest -R '_cpp$' # CPU-backend tests only (no GP Key CMake options: `ENABLE_CPU` (ON), `ENABLE_CUDA` (ON if nvcc found), `BUILD_TEST` (ON), `BUILD_UTEST` (ON), `ENABLE_BENCHMARK` (OFF), `CUDA_ARCH` ("native", passed verbatim to `CUDA_ARCHITECTURES` — no arch filtering exists), `ARCH_FLAGS` (CPU SIMD), `ENABLE_NVTX`, `ENABLE_DEBUG`, `TEMPLATE_DEPTH` (1000). -There is no lint/format gate. Format manually with `clang-format -i` using the repo-root `.clang-format` (LLVM base, 4-space indent, 120-char lines, `PointerAlignment: Right`). The merge gate is the full ctest suite building and passing across CI's compiler matrix (no `-Werror` anywhere, though the skills' PR checklists ask for warning-clean nvcc and clang builds). CI runs only on PRs to `main` (self-hosted runners): Linux amd64/arm64 with g++-13 and clang++-21 + CUDA 13.3; Windows with cl 14.44 + CUDA 13.0, cl 14.51 + CUDA 13.3, and clang-cl 14.51 + CUDA 13.3. +There is no lint/format gate. Format manually with `clang-format -i` using the repo-root `.clang-format` (LLVM base, 4-space indent, 120-char lines, `PointerAlignment: Right`). The merge gate is the full ctest suite building and passing across CI's compiler matrix (no `-Werror` anywhere, though the skills' PR checklists ask for warning-clean nvcc and clang builds). CI runs only on PRs to `main` (self-hosted runners): Linux amd64/arm64 with g++-13 and clang++-23 + CUDA 13.3; Windows with cl 14.44 + CUDA 13.0, cl 14.51 + CUDA 13.3, and clang-cl 14.51 + CUDA 13.3. ## Test infrastructure (no framework — no GTest/Catch2) diff --git a/CMakeLists.txt b/CMakeLists.txt index 6dea13c1..5b38cce0 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1,6 +1,8 @@ cmake_minimum_required(VERSION 3.28 FATAL_ERROR) - +set(CMAKE_CXX_SCAN_FOR_MODULES OFF) +set(CMAKE_EXPORT_COMPILE_COMMANDS ON) +set(CMAKE_COLOR_DIAGNOSTICS ON) set (PROJECT_VERSION_MAJOR 0) set (PROJECT_VERSION_MINOR 2) set (PROJECT_VERSION_REV 1) @@ -18,15 +20,43 @@ HOMEPAGE_URL "https://github.com/morousg/FusedKernelLibrary" ) #cuda is optional, but if it is found, it will be used option(ENABLE_CPU "Enable CPU support" ON) +option(ENABLE_CUDA "Enable CUDA support" ON) +option(ENABLE_HIP "Enable HIP support" OFF) + +if (ENABLE_CUDA AND ENABLE_HIP) + message(FATAL_ERROR "ENABLE_CUDA and ENABLE_HIP cannot be enabled at the same time.") +endif() + include(CheckLanguage) -check_language(CUDA) -if (CMAKE_CUDA_COMPILER) - option (ENABLE_CUDA "Enable CUDA support" ON) - if (${ENABLE_CUDA}) + +if (ENABLE_CUDA) + check_language(CUDA) + if (CMAKE_CUDA_COMPILER) include(cmake/cuda_init.cmake) + else() + message(WARNING "CUDA compiler not found. Disabling CUDA support") + set(ENABLE_CUDA OFF CACHE BOOL "Enable CUDA support" FORCE) + endif() +endif() + +if (ENABLE_HIP) + check_language(HIP) + if (CMAKE_HIP_COMPILER) + if(NOT DEFINED HIP_PATH) + if(NOT DEFINED ENV{HIP_PATH}) + set(HIP_PATH "/opt/rocm/hip" CACHE PATH "Path to which HIP has been installed") + else() + set(HIP_PATH $ENV{HIP_PATH} CACHE PATH "Path to which HIP has been installed") + endif() + endif() + set(CMAKE_MODULE_PATH "${HIP_PATH}/cmake" ${CMAKE_MODULE_PATH}) + set(CMAKE_CXX_SCAN_FOR_MODULES OFF) #clang with HIP does not come with clang-scan-deps + include(cmake/hip_init.cmake) + message("Using HIP compiler at " ${CMAKE_HIP_COMPILER}) + else() + message(WARNING "HIP compiler not found. Disabling HIP support") + set(ENABLE_HIP OFF CACHE BOOL "Enable HIP support" FORCE) endif() -else() - message(STATUS "CUDA compiler not found, CUDA support will be disabled.") endif() add_subdirectory(include) diff --git a/README.md b/README.md index c597669e..0d4fbf03 100644 --- a/README.md +++ b/README.md @@ -1,12 +1,12 @@ # Fused Kernel Library (FKL) Main branch -The Fused Kernel Library is a C++20 implementation of a methodology that allows to define a set of operations that need to be executed inside a kernel, in the same order that they are expressed. The library currently has CPU and CUDA backends, but other GPU language implemenations (such as HIP) are possible. +The Fused Kernel Library is a C++20 implementation of a methodology that allows to define a set of operations that need to be executed inside a kernel, in the same order that they are expressed. The library currently has CPU, CUDA and ROCM backends, but other GPU language implementations are possible. It automatically implements Vertical and Horizontal fusion, and also implements two new Fusion techniques, Backwards Vertical Fusion (akin to OpenCV Filters, but with an standard generic API), and Divergent Horizontal Fusion. This is the Main branch of the repository, where the FKL API can be modified, the minimum C++ version can be increased (currently it is C++20), and in general, we can break retro-compatibility with the goal of improving future code maintainability, code performance, and being able to fuse more types of algorithms. -If you are looking for a more stable branch, where new features can be added, but the FKL API is frozen to ensure your user code will continue to work with newer versions, check the LTS-C++17 branch. +If you are looking for a more stable branch, where new features can be added, but the FKL API is frozen to ensure your user code will continue to work with newer versions, check the LTS-C++17 branch (ROCM is not supported there). ## Reference paper and other publications This repository provides the official implementation of a kernel fusion methodology for GPU libraries, providing the mechanisms to perform automatic Vertical Fusion, Horizontal Fusion, Backwards Vertical Fusion and Divergent Horizontal Fusion. @@ -159,7 +159,7 @@ This has been tested before, by creating special compilers that generate the ass ## Closed source friendly -A company that has it's own CUDA kernels, and wants to start fusing them along with operations present in this library, can do so by shaping their kernels into a conformant FusedKernel Operation, that can be passed as a template parameter of one of the FKL InstantiableOperation structs. +A company that has it's own CUDA or HIP kernels, and wants to start fusing them along with operations present in this library, can do so by shaping their kernels into a conformant FusedKernel Operation, that can be passed as a template parameter of one of the FKL InstantiableOperation structs. With this strategy, they don't need to share any of their code. They just need to make their kernels fusionable. @@ -168,5 +168,9 @@ With this strategy, they don't need to share any of their code. They just need t This is an Apache 2.0 OpenSource project, currently with no funding. The 3 main contributors work on the project on their spare time. In the case of the main author, he can work on solving bugs or adding features during his work hours on Grup Mediapro S.L.U. The code added during that time will include Copyright Grup Mediapro S.L.U. As per Apache 2.0 we provide no guaratees, or free support. Nevertheless, the software has some testing on the following OS and compiler versions: -- Ubuntu 24.04 g++ 13 + CUDA 13.3 or clang21 + CUDA 13.3 for both amd64 and arm64 systems. +for CUDA builds (x86_64 and arm64) +- Ubuntu 24.04 g++ 13 + CUDA 13.3 or clang++-23 + CUDA 13.3 for both amd64 and arm64 systems. - Windows 11 Visual Studio 2022 + CUDA 13.0 or Visual Studio 2026 + CUDA 13.3 or clang-cl (VS2026) + CUDA 13.3 on amd64 systems. + +For ROCM builds (x86_64 only): +- Ubuntu 24.04 + ROCM 10.0. diff --git a/benchmarks/CMakeLists.txt b/benchmarks/CMakeLists.txt index 7aefe330..e332b266 100644 --- a/benchmarks/CMakeLists.txt +++ b/benchmarks/CMakeLists.txt @@ -15,6 +15,11 @@ function(add_cuda_to_benchmark TARGET_NAME) endif() endfunction() +function(add_hip_to_benchmark TARGET_NAME) + add_hip_to_target(${TARGET_NAME} "") + set_target_hip_arch_flags(${TARGET_NAME}) +endfunction() + function (add_generated_benchmark TARGET_NAME TEST_SOURCE EXTENSION DIR) set(TEST_GENERATED_SOURCE "${CMAKE_CURRENT_BINARY_DIR}/${TARGET_NAME}_${EXTENSION}/launcher.${EXTENSION}") #use the same name as the target ) @@ -27,13 +32,11 @@ function (add_generated_benchmark TARGET_NAME TEST_SOURCE EXTENSION DIR) if(${ENABLE_BENCHMARK}) target_compile_definitions(${TARGET_NAME_EXT} PRIVATE ENABLE_BENCHMARK) endif() + set_target_properties(${TARGET_NAME_EXT} PROPERTIES CXX_STANDARD 20 CXX_STANDARD_REQUIRED YES CXX_EXTENSIONS NO) - - #todo: add hip support - set_target_properties(${TARGET_NAME_EXT} PROPERTIES CXX_STANDARD 20 CXX_STANDARD_REQUIRED YES CXX_EXTENSIONS NO) target_include_directories(${TARGET_NAME_EXT} PRIVATE "${CMAKE_SOURCE_DIR}") target_include_directories(${TARGET_NAME_EXT} PRIVATE "${DIR}") - target_link_libraries(${TARGET_NAME_EXT} PRIVATE FKL::FKL) + target_link_libraries(${TARGET_NAME_EXT} PRIVATE FKL::FKL) if (MSVC) target_compile_options(${TARGET_NAME_EXT} PRIVATE $<$:/diagnostics:caret>) target_compile_options(${TARGET_NAME_EXT} PRIVATE $<$:/Zc:preprocessor>) @@ -67,6 +70,10 @@ function (discover_benchmark DIR) add_generated_benchmark("${TARGET_NAME}" "${benchmark_source}" "cu" "${DIR_NAME}") add_cuda_to_benchmark("${TARGET_NAME}_cu") endif() + if (CMAKE_HIP_COMPILER AND ENABLE_HIP) + add_generated_benchmark("${TARGET_NAME}" "${benchmark_source}" "hip" "${DIR_NAME}") + add_hip_to_benchmark("${TARGET_NAME}_hip") + endif() endforeach() endfunction() diff --git a/benchmarks/fkBenchmarksCommon.h b/benchmarks/fkBenchmarksCommon.h index 980ce361..d29fffa9 100644 --- a/benchmarks/fkBenchmarksCommon.h +++ b/benchmarks/fkBenchmarksCommon.h @@ -99,7 +99,7 @@ class TimeMarkerOne final : public TimeMarkerInterfaceOne { } }; -#if defined(__CUDACC__) || defined(__HIP__) +#if defined(__CUDACC__) template <> class TimeMarkerOne final : public TimeMarkerInterfaceOne { cudaEvent_t m_start, m_stop; @@ -130,7 +130,40 @@ class TimeMarkerOne final : public TimeMarkerInterfaceO return m_elapsedTime; } }; -#endif // defined(__CUDACC__) || defined(__HIP__) +#endif // defined(__CUDACC__) + +#if defined(__HIPCC__) +template <> +class TimeMarkerOne final : public TimeMarkerInterfaceOne { + hipEvent_t m_start, m_stop; + hipStream_t m_stream; + std::array m_elapsedTime; +public: + TimeMarkerOne(fk::Stream stream) : m_stream(stream) { + gpuErrchk(hipEventCreate(&m_start)); + gpuErrchk(hipEventCreate(&m_stop)); + m_elapsedTime.fill(0.f); + } + ~TimeMarkerOne() { + gpuErrchk(hipEventDestroy(m_start)); + gpuErrchk(hipEventDestroy(m_stop)); + } + void start() final { + gpuErrchk(hipEventRecord(m_start, m_stream)); + } + void stop(BenchmarkResultsNumbersOne& resF, const int& idx) final { + gpuErrchk(hipEventRecord(m_stop, m_stream)); + gpuErrchk(hipEventSynchronize(m_stop)); + gpuErrchk(hipEventElapsedTime(&m_elapsedTime[idx], m_start, m_stop)); + resF.fkElapsedTimeMax = resF.fkElapsedTimeMax < m_elapsedTime[idx] ? m_elapsedTime[idx] : resF.fkElapsedTimeMax; + resF.fkElapsedTimeMin = resF.fkElapsedTimeMin > m_elapsedTime[idx] ? m_elapsedTime[idx] : resF.fkElapsedTimeMin; + resF.fkElapsedTimeAcum += m_elapsedTime[idx]; + } + std::array getElapsedTime() const final { + return m_elapsedTime; + } +}; +#endif // defined(__HIPCC__) #if defined(NVRTC_ENABLED) template <> @@ -226,7 +259,7 @@ class TimeMarkerTwo final : public TimeMarkerInterfaceTwo { }; }; -#if defined(__CUDACC__) || defined(__HIP__) +#if defined(__CUDACC__) template <> class TimeMarkerTwo final : public TimeMarkerInterfaceTwo { cudaEvent_t m_start, m_stop; @@ -273,7 +306,56 @@ class TimeMarkerTwo final : public TimeMarkerInterfaceT return m_secondElapsedTime; }; }; -#endif +#endif // defined(__CUDACC__) + +#if defined(__HIPCC__) +template <> +class TimeMarkerTwo final : public TimeMarkerInterfaceTwo { + hipEvent_t m_start, m_stop; + hipStream_t m_stream; + std::array m_firstElapsedTime; + std::array m_secondElapsedTime; +public: + TimeMarkerTwo(fk::Stream stream) : m_stream(stream) { + gpuErrchk(hipEventCreate(&m_start)); + gpuErrchk(hipEventCreate(&m_stop)); + m_firstElapsedTime.fill(0.f); + m_secondElapsedTime.fill(0.f); + } + + ~TimeMarkerTwo() { + gpuErrchk(hipEventDestroy(m_start)); + gpuErrchk(hipEventDestroy(m_stop)); + } + + void startFirst() final { + gpuErrchk(hipEventRecord(m_start, m_stream)); + }; + void stopFirstStartSecond(BenchmarkResultsNumbersTwo& resF, const int& idx) final { + gpuErrchk(hipEventRecord(m_stop, m_stream)); + gpuErrchk(hipEventSynchronize(m_stop)); + gpuErrchk(hipEventElapsedTime(&m_firstElapsedTime[idx], m_start, m_stop)); + resF.firstElapsedTimeMax = resF.firstElapsedTimeMax < m_firstElapsedTime[idx] ? m_firstElapsedTime[idx] : resF.firstElapsedTimeMax; + resF.firstElapsedTimeMin = resF.firstElapsedTimeMin > m_firstElapsedTime[idx] ? m_firstElapsedTime[idx] : resF.firstElapsedTimeMin; + resF.firstElapsedTimeAcum += m_firstElapsedTime[idx]; + gpuErrchk(hipEventRecord(m_start, m_stream)); + } + void stopSecond(BenchmarkResultsNumbersTwo& resF, const int& idx) final { + gpuErrchk(hipEventRecord(m_stop, m_stream)); + gpuErrchk(hipEventSynchronize(m_stop)); + gpuErrchk(hipEventElapsedTime(&m_secondElapsedTime[idx], m_start, m_stop)); + resF.secondElapsedTimeMax = resF.secondElapsedTimeMax < m_secondElapsedTime[idx] ? m_secondElapsedTime[idx] : resF.secondElapsedTimeMax; + resF.secondElapsedTimeMin = resF.secondElapsedTimeMin > m_secondElapsedTime[idx] ? m_secondElapsedTime[idx] : resF.secondElapsedTimeMin; + resF.secondElapsedTimeAcum += m_secondElapsedTime[idx]; + }; + std::array getFirstElapsedTime() const final { + return m_firstElapsedTime; + }; + std::array getSecondElapsedTime() const final { + return m_secondElapsedTime; + }; +}; +#endif // defined(__HIPCC__) #if defined(NVRTC_ENABLED) template <> diff --git a/cmake/discover_tests.cmake b/cmake/discover_tests.cmake index f4160eb6..3324b4ba 100644 --- a/cmake/discover_tests.cmake +++ b/cmake/discover_tests.cmake @@ -15,46 +15,6 @@ function(add_cuda_to_test TARGET_NAME) endif() endfunction() -function (add_generated_test TARGET_NAME TEST_SOURCE EXTENSION DIR) - - set(TEST_GENERATED_SOURCE "${CMAKE_CURRENT_BINARY_DIR}/${TARGET_NAME}_${EXTENSION}/launcher.${EXTENSION}") #use the same name as the target ) - - configure_file(${CMAKE_SOURCE_DIR}/tests/launcher.in ${TEST_GENERATED_SOURCE} @ONLY) #replace variables in the test source file - set (TARGET_NAME_EXT "${TARGET_NAME}_${EXTENSION}") - #message(STATUS "Adding test: ${TARGET_NAME_EXT} from ${TEST_GENERATED_SOURCE} and ${TEST_SOURCE}") - add_executable(${TARGET_NAME_EXT} "${TEST_GENERATED_SOURCE};${TEST_SOURCE}" ) - target_sources(${TARGET_NAME_EXT} PRIVATE ${LAUNCH_SOURCES}) - - if(${ENABLE_BENCHMARK}) - target_compile_definitions(${TARGET_NAME_EXT} PRIVATE ENABLE_BENCHMARK) - endif() - - set_target_properties(${TARGET_NAME_EXT} PROPERTIES CXX_STANDARD 20 CXX_STANDARD_REQUIRED YES CXX_EXTENSIONS NO) - target_include_directories(${TARGET_NAME_EXT} PRIVATE "${CMAKE_SOURCE_DIR}") - target_include_directories(${TARGET_NAME_EXT} PRIVATE "${DIR}") - target_link_libraries(${TARGET_NAME_EXT} PRIVATE FKL::FKL) - target_link_libraries(${TARGET_NAME_EXT} PRIVATE CUDA::cuda_driver) - if (NVRTC_ENABLE) - target_link_libraries(${TARGET_NAME_EXT} PRIVATE ${NVRTC_LIBRARIES}) - target_compile_definitions(${TARGET_NAME_EXT} PRIVATE NVRTC_ENABLE) - if (MSVC) - if (NVRTC_STATIC_LINK) - set_target_properties(${TARGET_NAME_EXT} PROPERTIES MSVC_RUNTIME_LIBRARY "MultiThreaded$<$:Debug>") - endif() - endif() - endif() - if (MSVC) - target_compile_options(${TARGET_NAME_EXT} PRIVATE $<$:/diagnostics:caret>) - endif() - - add_optimization_flags(${TARGET_NAME_EXT}) - - add_test(NAME ${TARGET_NAME_EXT} COMMAND ${TARGET_NAME_EXT}) - cmake_path(SET path2 "${DIR}") - cmake_path(GET path2 FILENAME DIR_NAME) - set_property(TARGET "${TARGET_NAME_EXT}" PROPERTY FOLDER "tests/${EXTENSION}/${DIR_NAME}") - -endfunction() function (discover_tests DIR) file( diff --git a/cmake/hip_init.cmake b/cmake/hip_init.cmake new file mode 100644 index 00000000..22479dec --- /dev/null +++ b/cmake/hip_init.cmake @@ -0,0 +1,3 @@ +enable_language(HIP) +include(cmake/libs/hip/hip.cmake) +include(cmake/libs/hip/archs.cmake) diff --git a/cmake/libs/hip/archs.cmake b/cmake/libs/hip/archs.cmake new file mode 100644 index 00000000..45ebd6c4 --- /dev/null +++ b/cmake/libs/hip/archs.cmake @@ -0,0 +1,7 @@ +set(CMAKE_HIP_ARCHITECTURES OFF) + +set(HIP_ARCH "gfx90a" CACHE STRING "HIP architecture(s) to build (e.g. gfx90a, gfx1100)") + +function(set_target_hip_arch_flags TARGET_NAME) + set_target_properties(${TARGET_NAME} PROPERTIES HIP_ARCHITECTURES "${HIP_ARCH}") +endfunction() diff --git a/cmake/libs/hip/debug.cmake b/cmake/libs/hip/debug.cmake new file mode 100644 index 00000000..c0511ef3 --- /dev/null +++ b/cmake/libs/hip/debug.cmake @@ -0,0 +1,7 @@ +function(add_hip_debug_support_to_target TARGET_NAME) + target_compile_options(${TARGET_NAME} PRIVATE $<$,$>:-g>) +endfunction() + +function(add_hip_lineinfo_to_target TARGET_NAME) + target_compile_options(${TARGET_NAME} PRIVATE $<$:-gline-tables-only>) +endfunction() diff --git a/cmake/libs/hip/hip.cmake b/cmake/libs/hip/hip.cmake new file mode 100644 index 00000000..884c26ee --- /dev/null +++ b/cmake/libs/hip/hip.cmake @@ -0,0 +1,23 @@ +include(cmake/libs/hip/debug.cmake) +include(cmake/libs/hip/target_generation.cmake) +if (WIN32) + list(APPEND CMAKE_PREFIX_PATH "${HIP_PATH}") +else() +#linux + list(APPEND CMAKE_PREFIX_PATH "${HIP_PATH}") +endif() +find_package(HIP CONFIG REQUIRED) +option(ENABLE_HIP_LINE_INFO "Enable line info for HIP kernels compilation" ON) + +function(add_hip_to_target TARGET_NAME COMPONENTS) + set_default_hip_target_properties(${TARGET_NAME}) + + if (ENABLE_DEBUG) + add_hip_debug_support_to_target(${TARGET_NAME}) + endif() + if (ENABLE_HIP_LINE_INFO) + add_hip_lineinfo_to_target(${TARGET_NAME}) + endif() + target_link_libraries(${TARGET_NAME} PRIVATE hip::device) + +endfunction() diff --git a/cmake/libs/hip/target_generation.cmake b/cmake/libs/hip/target_generation.cmake new file mode 100644 index 00000000..b3548570 --- /dev/null +++ b/cmake/libs/hip/target_generation.cmake @@ -0,0 +1,9 @@ +function(set_default_hip_target_properties TARGET_NAME) + set_target_properties(${TARGET_NAME} PROPERTIES HIP_STANDARD_REQUIRED ON HIP_STANDARD 20 HIP_RUNTIME_LIBRARY Shared) + + set_target_hip_arch_flags(${TARGET_NAME}) + + if (NOT(${TEMPLATE_DEPTH} STREQUAL "default")) + target_compile_options(${TARGET_NAME} PRIVATE $<$:-ftemplate-depth=${TEMPLATE_DEPTH}>) + endif() +endfunction() diff --git a/cmake/tests/add_generated_test.cmake b/cmake/tests/add_generated_test.cmake index f1da5730..07ef78da 100644 --- a/cmake/tests/add_generated_test.cmake +++ b/cmake/tests/add_generated_test.cmake @@ -3,6 +3,7 @@ if (WIN32) list(APPEND LAUNCH_SOURCES "${CMAKE_SOURCE_DIR}/utf8cp.manifest") #for utf8 codepage endif() include (cmake/generators/export_header.cmake) + function(add_cuda_to_test TARGET_NAME) add_cuda_to_target(${TARGET_NAME} "") set_target_cuda_arch_flags(${TARGET_NAME}) @@ -15,6 +16,11 @@ function(add_cuda_to_test TARGET_NAME) endif() endfunction() +function(add_hip_to_test TARGET_NAME) + add_hip_to_target(${TARGET_NAME} "") + set_target_hip_arch_flags(${TARGET_NAME}) +endfunction() + function(configure_test_target_flags TARGET_NAME TEST_SOURCE DIR) set(TEST_GENERATED_SOURCE "${CMAKE_CURRENT_BINARY_DIR}/${TARGET_NAME}_${EXTENSION}/launcher.${EXTENSION}") #use the same name as the target ) @@ -25,9 +31,11 @@ function(configure_test_target_flags TARGET_NAME TEST_SOURCE DIR) target_compile_definitions(${TARGET_NAME} PRIVATE ENABLE_BENCHMARK) endif() set_target_properties(${TARGET_NAME} PROPERTIES CXX_STANDARD 20 CXX_STANDARD_REQUIRED YES CXX_EXTENSIONS NO) - target_include_directories(${TARGET_NAME} PUBLIC "${CMAKE_SOURCE_DIR}") - target_include_directories(${TARGET_NAME} PUBLIC "${DIR}") - target_link_libraries(${TARGET_NAME} PUBLIC FKL::FKL) + + + target_include_directories(${TARGET_NAME} PUBLIC "${CMAKE_SOURCE_DIR}") + target_include_directories(${TARGET_NAME} PUBLIC "${DIR}") + target_link_libraries(${TARGET_NAME} PUBLIC FKL::FKL) if (MSVC) target_compile_options(${TARGET_NAME} PRIVATE $<$:/diagnostics:caret>) target_compile_options(${TARGET_NAME} PRIVATE $<$:/bigobj>) @@ -45,9 +53,11 @@ endfunction() function (set_ide_target_folder TARGET_NAME DIR_PARENT_PATH EXTENSION) if (${EXTENSION} STREQUAL "cu") set(FKL_BACKEND "cuda") - elseif(${EXTENSION} STREQUAL "cpp") - set(FKL_BACKEND "cpu") - else() + elseif(${EXTENSION} STREQUAL "hip") + set(FKL_BACKEND "hip") + elseif(${EXTENSION} STREQUAL "cpp") + set(FKL_BACKEND "cpu") + else() message(FATAL_ERROR "Unknown extension: ${EXTENSION}") endif() set_property(TARGET "${TARGET_NAME}" PROPERTY FOLDER "${DIR_PARENT_PATH}/${FKL_BACKEND}/") @@ -55,8 +65,10 @@ endfunction() function (add_generated_test_stub TARGET_NAME_EXT TEST_SOURCE DIR) - set(TEST_GENERATED_SOURCE "${CMAKE_CURRENT_BINARY_DIR}/${TARGET_NAME}_${EXTENSION}/launcher.${EXTENSION}") #use the same name as the target ) + set(TEST_GENERATED_SOURCE "${CMAKE_CURRENT_BINARY_DIR}/${TARGET_NAME}_${EXTENSION}/launcher.${EXTENSION}") #use the same name as the target ) + configure_file(${CMAKE_SOURCE_DIR}/tests/launcher.in ${TEST_GENERATED_SOURCE} @ONLY) #replace variables in the test source file + add_executable(${TARGET_NAME_EXT} "${TEST_GENERATED_SOURCE};${TEST_SOURCE}" ) configure_test_target_flags("${TARGET_NAME_EXT}" "${TEST_SOURCE}" "${DIR}") target_sources(${TARGET_NAME_EXT} PRIVATE ${LAUNCH_SOURCES}) @@ -65,7 +77,15 @@ function (add_generated_test_stub TARGET_NAME_EXT TEST_SOURCE DIR) cmake_path(GET path2 FILENAME DIR_NAME) cmake_path(GET path2 PARENT_PATH DIR_PARENT_PATH) string(REPLACE " " "" DIR_PARENT_PATH "${DIR_PARENT_PATH}") #for windows - + if (${EXTENSION} STREQUAL "cu") + set(FKL_BACKEND "cuda") + set_source_files_properties("${TEST_GENERATED_SOURCE}" PROPERTIES LANGUAGE CUDA) #set the language of the source file to CUDA + elseif (${EXTENSION} STREQUAL "hip") + set(FKL_BACKEND "hip") + set_source_files_properties("${TEST_GENERATED_SOURCE}" PROPERTIES LANGUAGE HIP) #set the language of the source file to HIP + elseif (${EXTENSION} STREQUAL "cpp") + set_source_files_properties("${TEST_GENERATED_SOURCE}" PROPERTIES LANGUAGE CXX) #set the language of the source file to CUDA + endif() set_ide_target_folder("${TARGET_NAME_EXT}" "${DIR_PARENT_PATH}" "${EXTENSION}") endfunction() diff --git a/cmake/tests/add_shared_test_libs.cmake b/cmake/tests/add_shared_test_libs.cmake index 6ed34494..c438685b 100644 --- a/cmake/tests/add_shared_test_libs.cmake +++ b/cmake/tests/add_shared_test_libs.cmake @@ -34,6 +34,8 @@ function (add_shared_test_lib TARGET_BASE_NAME DIR EXTENSION FUNDAMENTAL_TYPE) add_shared_target("${TARGET_BASE_NAME}" "${EXTENSION}" "${FUNDAMENTAL_TYPE}" "${DIR}") if ("${EXTENSION}" STREQUAL "cu") add_cuda_to_test("${TARGET_NAME}_${EXTENSION}") + elseif ("${EXTENSION}" STREQUAL "hip") + add_hip_to_test("${TARGET_NAME}_${EXTENSION}") endif() endfunction() \ No newline at end of file diff --git a/cmake/tests/discover_tests.cmake b/cmake/tests/discover_tests.cmake index 57416626..f7c2d589 100644 --- a/cmake/tests/discover_tests.cmake +++ b/cmake/tests/discover_tests.cmake @@ -31,7 +31,14 @@ function (discover_tests DIR) add_generated_test("${TARGET_NAME}" "${TEST_SOURCE}" "cu" "${DIR_RELATIVE_PATH}") add_cuda_to_test("${TARGET_NAME}_cu") endif() - endif() + endif() + + if (CMAKE_HIP_COMPILER AND ENABLE_HIP) + if (${POS_ONLY_CPU} EQUAL -1 AND ${POS_ONLY_CU} EQUAL -1) + add_generated_test("${TARGET_NAME}" "${TEST_SOURCE}" "hip" "${DIR_RELATIVE_PATH}") + add_hip_to_test("${TARGET_NAME}_hip") + endif() + endif() endforeach() endfunction() \ No newline at end of file diff --git a/include/fused_kernel/algorithms/basic_ops/logical.h b/include/fused_kernel/algorithms/basic_ops/logical.h index a6e07d83..d5dda813 100644 --- a/include/fused_kernel/algorithms/basic_ops/logical.h +++ b/include/fused_kernel/algorithms/basic_ops/logical.h @@ -100,18 +100,21 @@ namespace fk { } template > FK_HOST_DEVICE_FUSE std::enable_if_t exec(const InputType input) { - const auto result = get<0>(input) == get<1>(input); - return result.x && result.y; + return get<0>(input).x == get<1>(input).x && + get<0>(input).y == get<1>(input).y; } template > FK_HOST_DEVICE_FUSE std::enable_if_t exec(const InputType input) { - const auto result = get<0>(input) == get<1>(input); - return result.x && result.y && result.z; + return get<0>(input).x == get<1>(input).x && + get<0>(input).y == get<1>(input).y && + get<0>(input).z == get<1>(input).z; } template > FK_HOST_DEVICE_FUSE std::enable_if_t exec(const InputType input) { - const auto result = get<0>(input) == get<1>(input); - return result.x && result.y && result.z && result.w; + return get<0>(input).x == get<1>(input).x && + get<0>(input).y == get<1>(input).y && + get<0>(input).z == get<1>(input).z && + get<0>(input).w == get<1>(input).w; } }; } //namespace fk diff --git a/include/fused_kernel/algorithms/image_processing/color_conversion.h b/include/fused_kernel/algorithms/image_processing/color_conversion.h index 27b429de..daa69026 100644 --- a/include/fused_kernel/algorithms/image_processing/color_conversion.h +++ b/include/fused_kernel/algorithms/image_processing/color_conversion.h @@ -114,7 +114,10 @@ namespace fk { using Parent = UnaryOperation>, O, DenormalizePixel>; DECLARE_UNARY_PARENT FK_HOST_DEVICE_FUSE OutputType exec(const InputType input) { - constexpr auto maxDepth = maxDepthValue; + // Keep maxDepth as float: multiplying a float vector by an integer scalar (e.g. uint + // for p8bit) can resolve to HIP's native vector-type operator, which narrows the + // vector to the scalar's type and performs the multiplication in integer arithmetic. + constexpr float maxDepth = static_cast(maxDepthValue); return cxp::cast::f(input * maxDepth); } }; @@ -128,7 +131,12 @@ namespace fk { using Parent = UnaryOperation>, NormalizePixel>; DECLARE_UNARY_PARENT FK_HOST_DEVICE_FUSE OutputType exec(const InputType input) { - return input / static_cast(maxDepthValue); + // Cast to float vector first: dividing an integer vector (e.g. uchar3) directly by a + // float scalar can be resolved to the backend's native vector-type operator (HIP in + // particular), which narrows the scalar to the vector's base type and performs the + // division in integer arithmetic, truncating the result to 0 for most inputs. + const OutputType floatInput = cxp::cast::f(input); + return floatInput / static_cast(maxDepthValue); } }; diff --git a/include/fused_kernel/algorithms/image_processing/deinterlace.h b/include/fused_kernel/algorithms/image_processing/deinterlace.h index fa1bb7a0..363bd43d 100644 --- a/include/fused_kernel/algorithms/image_processing/deinterlace.h +++ b/include/fused_kernel/algorithms/image_processing/deinterlace.h @@ -80,11 +80,17 @@ namespace fk { using ReadOperation = typename BackIOp::Operation; // Read current pixel - const auto current = ReadOperation::exec(thread, backIOp); + const auto current = cxp::cast::f(ReadOperation::exec(thread, backIOp)); - if (thread.y > 0) { - const auto above = ReadOperation::exec(Point{thread.x, thread.y - 1, thread.z}, backIOp); - return (current + above + 1) * 0.5f; + const auto height = ReadOperation::num_elems_y(Point{0,0,0}, backIOp); + if (thread.y > 0 && thread.y + 1 < height) { + const auto above = cxp::cast::f( + ReadOperation::exec(Point{thread.x, thread.y - 1, thread.z}, backIOp)); + + const auto below = cxp::cast::f( + ReadOperation::exec(Point{thread.x, thread.y + 1, thread.z}, backIOp)); + + return (current + above + below + 1.f) * 0.25f; } else { return current * 1.f; } @@ -94,11 +100,15 @@ namespace fk { using ReadOperation = typename BackIOp::Operation; if (interpolate) { // We average the above pixel with the below pixel - const auto above = ReadOperation::exec(Point{thread.x, thread.y - 1, thread.z}, backIOp); - const auto below = ReadOperation::exec(Point{thread.x, thread.y + 1, thread.z}, backIOp); - return (above + below + 1) * 0.5f; + const auto above = cxp::cast::f( + ReadOperation::exec(Point{thread.x, thread.y - 1, thread.z}, backIOp)); + + const auto below = cxp::cast::f( + ReadOperation::exec(Point{thread.x, thread.y + 1, thread.z}, backIOp)); + + return (above + below + 1.f) * 0.5f; } else { - return ReadOperation::exec(thread, backIOp) * 1.f; + return cxp::cast::f(ReadOperation::exec(thread, backIOp)); } } @@ -109,9 +119,11 @@ namespace fk { // If useEvenLines is true, we interpolate on odd lines, otherwise we interpolate the even lines // useEvenLines = true, we interpolate if thread.y is odd and not the last line // useEvenLines = false, we interpolate if thread.y is even and not the first line - const bool interpolate = params.useEvenLines ? - !cxp::is_even::f(thread.y) && thread.y != ReadOperation::num_elems_y(Point{0,0,0}, backIOp) - 1 - : cxp::is_even::f(thread.y) && thread.y != 0; + const bool interpolate = params.useEvenLines + ? !cxp::is_even::f(thread.y) + && thread.y != ReadOperation::num_elems_y(Point{0,0,0}, backIOp) - 1 + : cxp::is_even::f(thread.y) && thread.y != 0 + && thread.y != ReadOperation::num_elems_y(Point{0,0,0}, backIOp) - 1; return execInterLinearGetPixel(thread, backIOp, interpolate); } diff --git a/include/fused_kernel/algorithms/image_processing/image.h b/include/fused_kernel/algorithms/image_processing/image.h index 0f3fe0df..d8fe0717 100644 --- a/include/fused_kernel/algorithms/image_processing/image.h +++ b/include/fused_kernel/algorithms/image_processing/image.h @@ -82,6 +82,21 @@ namespace fk { inline void download(Stream_& stream) { data.download(stream); } +#elif defined(__HIPCC__) + inline void uploadTo(Image& other, hipStream_t stream = 0) { + data.uploadTo(other.data, stream); + } + + inline void downloadTo(Image& other, hipStream_t stream = 0) { + data.downloadTo(other.data, stream); + } + + inline void upload(Stream_& stream) { + data.upload(stream); + } + inline void download(Stream_& stream) { + data.download(stream); + } #else inline void upload(Stream& stream) {} inline void download(Stream& stream) {} diff --git a/include/fused_kernel/algorithms/image_processing/interpolation.h b/include/fused_kernel/algorithms/image_processing/interpolation.h index edb76ff3..e873a6af 100644 --- a/include/fused_kernel/algorithms/image_processing/interpolation.h +++ b/include/fused_kernel/algorithms/image_processing/interpolation.h @@ -71,7 +71,7 @@ namespace fk { const float src_x = input.x; const float src_y = input.y; -#ifdef __CUDA_ARCH__ +#if defined(__CUDA_ARCH__) || defined(__HIP_ARCH__) const int x1 = __float2int_rd(src_x); const int y1 = __float2int_rd(src_y); #else @@ -91,10 +91,12 @@ namespace fk { {x2_read, y2_read, 0} }; // Read the 4 pixels from backIOp Read or ReadBack Operation - const auto src_reg0x0 = BackIOp::Operation::exec(readPoints._0x0, backIOp); - const auto src_reg1x0 = BackIOp::Operation::exec(readPoints._1x0, backIOp); - const auto src_reg0x1 = BackIOp::Operation::exec(readPoints._0x1, backIOp); - const auto src_reg1x1 = BackIOp::Operation::exec(readPoints._1x1, backIOp); + + const auto src_reg0x0 = cxp::cast::f(BackIOp::Operation::exec(readPoints._0x0, backIOp)); + const auto src_reg1x0 = cxp::cast::f(BackIOp::Operation::exec(readPoints._1x0, backIOp)); + const auto src_reg0x1 = cxp::cast::f(BackIOp::Operation::exec(readPoints._0x1, backIOp)); + const auto src_reg1x1 = cxp::cast::f(BackIOp::Operation::exec(readPoints._1x1, backIOp)); + // Compute the interpolated pixel and return it return (src_reg0x0 * ((x2 - src_x) * (y2 - src_y))) + diff --git a/include/fused_kernel/core/data/ptr_nd.h b/include/fused_kernel/core/data/ptr_nd.h index e562e637..aa373a6f 100644 --- a/include/fused_kernel/core/data/ptr_nd.h +++ b/include/fused_kernel/core/data/ptr_nd.h @@ -27,7 +27,7 @@ namespace fk { enum class MemType { Device, Host, HostPinned, DeviceAndPinned }; -#if defined(__NVCC__) +#if defined(__NVCC__) || defined(__HIPCC__) constexpr MemType defaultMemType = MemType::DeviceAndPinned; #else constexpr MemType defaultMemType = MemType::Host; @@ -48,7 +48,11 @@ namespace fk { if (ptr_a.dims.pitch == 0) { ptr_a.dims.pitch = sizeof(T) * ptr_a.dims.width; } +#if defined(__HIPCC__) + gpuErrchk(hipMalloc(&ptr_a.data, ptr_a.dims.pitch)); +#else gpuErrchk(cudaMalloc(&ptr_a.data, ptr_a.dims.pitch)); +#endif } FK_HOST_FUSE void h_malloc_init(PtrDims& dims) { if (dims.pitch == 0) { @@ -68,10 +72,18 @@ namespace fk { FK_HOST_FUSE void d_malloc(RawPtr& ptr_a) { if (ptr_a.dims.pitch == 0) { size_t pitch; +#if defined(__HIPCC__) + gpuErrchk(hipMallocPitch(&ptr_a.data, &pitch, sizeof(T) * ptr_a.dims.width, ptr_a.dims.height)); +#else gpuErrchk(cudaMallocPitch(&ptr_a.data, &pitch, sizeof(T) * ptr_a.dims.width, ptr_a.dims.height)); +#endif ptr_a.dims.pitch = static_cast(pitch); } else { +#if defined(__HIPCC__) + gpuErrchk(hipMalloc(&ptr_a.data, PtrImpl::sizeInBytes(ptr_a.dims))); +#else gpuErrchk(cudaMalloc(&ptr_a.data, PtrImpl::sizeInBytes(ptr_a.dims))); +#endif } } FK_HOST_FUSE void h_malloc_init(PtrDims& dims) { @@ -93,7 +105,11 @@ namespace fk { if (ptr_a.dims.pitch == 0) { ptr_a.dims.pitch = sizeof(T) * ptr_a.dims.width; } +#if defined(__HIPCC__) + gpuErrchk(hipMalloc(&ptr_a.data, PtrImpl::sizeInBytes(ptr_a.dims))); +#else gpuErrchk(cudaMalloc(&ptr_a.data, PtrImpl::sizeInBytes(ptr_a.dims))); +#endif ptr_a.dims.plane_pitch = ptr_a.dims.pitch * ptr_a.dims.height; } FK_HOST_FUSE void h_malloc_init(PtrDims& dims) { @@ -114,7 +130,11 @@ namespace fk { } FK_HOST_FUSE void d_malloc(RawPtr& ptr_a) { PtrImpl::h_malloc_init(ptr_a.dims); +#if defined(__HIPCC__) + gpuErrchk(hipMalloc(&ptr_a.data, PtrImpl::sizeInBytes(ptr_a.dims))); +#else gpuErrchk(cudaMalloc(&ptr_a.data, PtrImpl::sizeInBytes(ptr_a.dims))); +#endif } FK_HOST_FUSE void h_malloc_init(PtrDims& dims) { if (dims.pitch == 0) { @@ -152,16 +172,25 @@ namespace fk { } inline constexpr void allocDevice() { - #if defined(__NVCC__) + #if defined(__NVCC__) || defined(__HIPCC__) int currentDevice; +#if defined(__HIPCC__) + gpuErrchk(hipGetDevice(¤tDevice)); + gpuErrchk(hipSetDevice(deviceID)); +#else gpuErrchk(cudaGetDevice(¤tDevice)); gpuErrchk(cudaSetDevice(deviceID)); +#endif PtrImpl::d_malloc(ptr_a); if (currentDevice != deviceID) { +#if defined(__HIPCC__) + gpuErrchk(hipSetDevice(currentDevice)); +#else gpuErrchk(cudaSetDevice(currentDevice)); +#endif } #else - throw std::runtime_error("Device allocation not supported in non-CUDA compilation."); + throw std::runtime_error("Device allocation not supported in non-GPU compilation."); #endif } @@ -171,33 +200,59 @@ namespace fk { } inline constexpr void allocHostPinned() { - #if defined(__NVCC__) + #if defined(__NVCC__) || defined(__HIPCC__) int currentDevice; +#if defined(__HIPCC__) + gpuErrchk(hipGetDevice(¤tDevice)); + gpuErrchk(hipSetDevice(deviceID)); +#else gpuErrchk(cudaGetDevice(¤tDevice)); gpuErrchk(cudaSetDevice(deviceID)); - PtrImpl::h_malloc_init(ptr_a.dims); +#endif + PtrImpl::h_malloc_init(ptr_a.dims); +#if defined(__HIPCC__) + gpuErrchk(hipHostMalloc(&ptr_a.data, PtrImpl::sizeInBytes(ptr_a.dims))); +#else gpuErrchk(cudaMallocHost(&ptr_a.data, PtrImpl::sizeInBytes(ptr_a.dims))); +#endif if (currentDevice != deviceID) { +#if defined(__HIPCC__) + gpuErrchk(hipSetDevice(currentDevice)); +#else gpuErrchk(cudaSetDevice(currentDevice)); +#endif } #else - throw std::runtime_error("Host pinned allocation not supported in non-CUDA compilation."); + throw std::runtime_error("Host pinned allocation not supported in non-GPU compilation."); #endif } inline constexpr void allocDeviceAndPinned() { - #if defined(__NVCC__) + #if defined(__NVCC__) || defined(__HIPCC__) int currentDevice; +#if defined(__HIPCC__) + gpuErrchk(hipGetDevice(¤tDevice)); + gpuErrchk(hipSetDevice(deviceID)); +#else gpuErrchk(cudaGetDevice(¤tDevice)); gpuErrchk(cudaSetDevice(deviceID)); +#endif PtrImpl::d_malloc(ptr_a); PtrImpl::h_malloc_init(ptr_pinned.dims); +#if defined(__HIPCC__) + gpuErrchk(hipHostMalloc(&ptr_pinned.data, PtrImpl::sizeInBytes(ptr_pinned.dims))); +#else gpuErrchk(cudaMallocHost(&ptr_pinned.data, PtrImpl::sizeInBytes(ptr_pinned.dims))); +#endif if (currentDevice != deviceID) { +#if defined(__HIPCC__) + gpuErrchk(hipSetDevice(currentDevice)); +#else gpuErrchk(cudaSetDevice(currentDevice)); +#endif } #else - throw std::runtime_error("Host pinned and Device allocations not supported in non-CUDA compilation."); + throw std::runtime_error("Host pinned and Device allocations not supported in non-GPU compilation."); #endif } @@ -209,10 +264,14 @@ namespace fk { switch (type) { case MemType::Device: { - #if defined(__NVCC__) + #if defined(__NVCC__) || defined(__HIPCC__) +#if defined(__HIPCC__) + gpuErrchk(hipFree(ref->ptr)); +#else gpuErrchk(cudaFree(ref->ptr)); +#endif #else - throw std::runtime_error("Device memory deallocation not supported in non-CUDA compilation."); + throw std::runtime_error("Device memory deallocation not supported in non-GPU compilation."); #endif break; } @@ -223,20 +282,29 @@ namespace fk { } case MemType::HostPinned: { - #if defined(__NVCC__) + #if defined(__NVCC__) || defined(__HIPCC__) +#if defined(__HIPCC__) + gpuErrchk(hipHostFree(ref->ptr)); +#else gpuErrchk(cudaFreeHost(ref->ptr)); +#endif #else - throw std::runtime_error("Host pinned memory deallocation not supported in non-CUDA compilation."); + throw std::runtime_error("Host pinned memory deallocation not supported in non-GPU compilation."); #endif break; } case MemType::DeviceAndPinned: { -#if defined(__NVCC__) +#if defined(__NVCC__) || defined(__HIPCC__) +#if defined(__HIPCC__) + gpuErrchk(hipFree(ref->ptr)); + gpuErrchk(hipHostFree(ref->pinnedPtr)); +#else gpuErrchk(cudaFree(ref->ptr)); gpuErrchk(cudaFreeHost(ref->pinnedPtr)); +#endif #else - throw std::runtime_error("Device and Host pinned memory deallocation not supported in non-CUDA compilation."); + throw std::runtime_error("Device and Host pinned memory deallocation not supported in non-GPU compilation."); #endif break; } @@ -249,6 +317,7 @@ namespace fk { } } +#if defined(__NVCC__) || defined(__HIPCC__) #if defined(__NVCC__) inline void copy(const RawPtr& thisPtr, RawPtr& other, const cudaMemcpyKind& kind, cudaStream_t stream = 0) const { @@ -267,6 +336,25 @@ namespace fk { } } } +#else + inline void copy(const RawPtr& thisPtr, RawPtr& other, const hipMemcpyKind& kind, + hipStream_t stream = 0) const { + if ((other.dims.pitch == other.dims.width * sizeof(T)) && (thisPtr.dims.pitch == thisPtr.dims.width * sizeof(T))) { + if (sizeInBytes() != PtrImpl::sizeInBytes(other.dims)) { + throw std::runtime_error("Size mismatch in upload."); + } + const size_t totalBytes = sizeInBytes(); + gpuErrchk(hipMemcpyAsync(other.data, thisPtr.data, totalBytes, kind, stream)); + } else { + if constexpr (D > ND::_2D || D == ND::_1D) { + throw std::runtime_error("Padding only supported in 2D pointers"); + } else { + gpuErrchk(hipMemcpy2DAsync(other.data, other.dims.pitch, thisPtr.data, thisPtr.dims.pitch, + thisPtr.dims.width * sizeof(T), thisPtr.dims.height, kind, stream)); + } + } + } +#endif #endif public: @@ -516,7 +604,6 @@ namespace fk { throw std::runtime_error("Download can only copy from Device pointers."); } } - inline void upload(Stream_& stream) { if (type == MemType::DeviceAndPinned) { constexpr cudaMemcpyKind kind = cudaMemcpyHostToDevice; @@ -529,6 +616,55 @@ namespace fk { copy(ptr_a, ptr_pinned, kind, stream); } } +#elif defined(__HIPCC__) + inline void uploadTo(Ptr& other, hipStream_t stream = 0) { + constexpr hipMemcpyKind kind = hipMemcpyHostToDevice; + constexpr MemType otherExpectedMemType1 = MemType::Device; + constexpr MemType otherExpectedMemType2 = MemType::DeviceAndPinned; + constexpr MemType thisExpectedMemType1 = MemType::Host; + constexpr MemType thisExpectedMemType2 = MemType::HostPinned; + if (type == thisExpectedMemType1 || type == thisExpectedMemType2) { + if (other.getMemType() == otherExpectedMemType1 || other.getMemType() == otherExpectedMemType2) { + auto dstRawPtr = other.ptr(); + copy(ptr_a, dstRawPtr, kind, stream); + } else { + throw std::runtime_error("Upload can only copy to Device pointers"); + } + } else { + throw std::runtime_error("Upload can only copy from Host or HostPinned pointers."); + } + } + + inline void downloadTo(Ptr& other, hipStream_t stream = 0) { + constexpr hipMemcpyKind kind = hipMemcpyDeviceToHost; + constexpr MemType otherExpectedMemType1 = MemType::Host; + constexpr MemType otherExpectedMemType2 = MemType::HostPinned; + constexpr MemType thisExpectedMemType1 = MemType::Device; + constexpr MemType thisExpectedMemType2 = MemType::DeviceAndPinned; + if (type == thisExpectedMemType1 || type == thisExpectedMemType2) { + if (other.getMemType() == otherExpectedMemType1 || other.getMemType() == otherExpectedMemType2) { + auto dstRawPtr = other.ptr(); + copy(ptr_a, dstRawPtr, kind, stream); + } else { + throw std::runtime_error("Download can only copy to Host or HostPinned pointers."); + } + } else { + throw std::runtime_error("Download can only copy from Device pointers."); + } + } + + inline void upload(Stream_& stream) { + if (type == MemType::DeviceAndPinned) { + constexpr hipMemcpyKind kind = hipMemcpyHostToDevice; + copy(ptr_pinned, ptr_a, kind, stream); + } + } + inline void download(Stream_& stream) { + if (type == MemType::DeviceAndPinned) { + constexpr hipMemcpyKind kind = hipMemcpyDeviceToHost; + copy(ptr_a, ptr_pinned, kind, stream); + } + } #else inline void upload(Stream& stream) {} inline void download(Stream& stream) {} diff --git a/include/fused_kernel/core/data/ptr_utils.h b/include/fused_kernel/core/data/ptr_utils.h index 56dc1f55..016e9c77 100644 --- a/include/fused_kernel/core/data/ptr_utils.h +++ b/include/fused_kernel/core/data/ptr_utils.h @@ -23,17 +23,19 @@ namespace fk { template inline void setTo(const T& value, Ptr& outputPtr, Stream_& stream) { RawPtr output = outputPtr.ptr(); -#if defined(__NVCC__) - if constexpr (PA == ParArch::GPU_NVIDIA) { +#if defined(__NVCC__) || defined(__HIPCC__) + if constexpr (PA == ParArch::GPU_NVIDIA || PA == ParArch::GPU_AMD) { if (outputPtr.getMemType() == MemType::Device || outputPtr.getMemType() == MemType::DeviceAndPinned) { - Executor>::executeOperations(stream, ReadSet::build(value, outputPtr.dims()), PerThreadWrite::build(output)); + Executor>::executeOperations(stream, ReadSet::build(value, outputPtr.dims()), PerThreadWrite::build(output)); if (outputPtr.getMemType() == MemType::DeviceAndPinned) { Stream_ cpuStream; Executor>::executeOperations(cpuStream, ReadSet::build(value, outputPtr.dims()), PerThreadWrite::build(outputPtr.ptrPinned())); } } else { - Executor>::executeOperations(stream, ReadSet::build(value, outputPtr.dims()), PerThreadWrite::build(output)); + Stream_ cpuStream; + Executor>::executeOperations( + cpuStream, ReadSet::build(value, outputPtr.dims()), PerThreadWrite::build(output)); } } else { diff --git a/include/fused_kernel/core/data/vector_types.h b/include/fused_kernel/core/data/vector_types.h index 31fa2715..d8b87b53 100644 --- a/include/fused_kernel/core/data/vector_types.h +++ b/include/fused_kernel/core/data/vector_types.h @@ -259,7 +259,9 @@ namespace fk { }; } // namespace fk -#if defined(__VECTOR_TYPES_H__) || defined(__CUDACC__) || defined(__NVCC__) +#if defined(__HIPCC__) +#include +#elif defined(__VECTOR_TYPES_H__) || defined(__CUDACC__) || defined(__NVCC__) #include #else using char1 = fk::Char1; diff --git a/include/fused_kernel/core/execution_model/data_parallel_patterns.h b/include/fused_kernel/core/execution_model/data_parallel_patterns.h index 55241a4f..7a7efa7c 100644 --- a/include/fused_kernel/core/execution_model/data_parallel_patterns.h +++ b/include/fused_kernel/core/execution_model/data_parallel_patterns.h @@ -68,6 +68,7 @@ namespace fk { // namespace FusedKernel template struct TransformDPPBase { friend struct TransformDPP; // Allow TransformDPP to access private members + friend struct TransformDPP; // Allow TransformDPP to access private members friend struct TransformDPP; // Allow TransformDPPBase to access private members private: using Details = DPPDetails; @@ -203,14 +204,17 @@ namespace fk { // namespace FusedKernel } }; -#if defined(__NVCC__) - template - struct TransformDPP, void>> { +#if defined(__NVCC__) || defined(__HIPCC__) + template + struct TransformDPP && + (PA == ParArch::GPU_NVIDIA || PA == ParArch::GPU_AMD), + void>> { private: using Parent = TransformDPPBase; using Details = DPPDetails; public: - static constexpr ParArch PAR_ARCH = ParArch::GPU_NVIDIA; + static constexpr ParArch PAR_ARCH = PA; template FK_HOST_DEVICE_FUSE ActiveThreads getActiveThreads(const Details& details, const FirstIOp& iOp) { @@ -231,11 +235,9 @@ namespace fk { // namespace FusedKernel template FK_DEVICE_FUSE void exec(const Details& details, const IOps&... iOps) { - const cg::thread_block g = cg::this_thread_block(); - - const int x = (g.dim_threads().x * g.group_index().x) + g.thread_index().x; - const int y = (g.dim_threads().y * g.group_index().y) + g.thread_index().y; - const int z = g.group_index().z; // So far we only consider the option of using the z dimension to specify n (x*y) thread planes + const int x = (blockDim.x * blockIdx.x) + threadIdx.x; + const int y = (blockDim.y * blockIdx.y) + threadIdx.y; + const int z = blockIdx.z; const Point thread{ x, y, z }; exec_thread(thread, details, iOps...); @@ -289,6 +291,7 @@ namespace fk { // namespace FusedKernel template struct DivergentBatchTransformDPPBase { friend struct DivergentBatchTransformDPP; // Allow DivergentBatchTransformDPP to access private members + friend struct DivergentBatchTransformDPP; // Allow DivergentBatchTransformDPP to access private members friend struct DivergentBatchTransformDPP; // Allow DivergentBatchTransformDPPBase to access private members private: template @@ -337,12 +340,15 @@ namespace fk { // namespace FusedKernel template <> struct DivergentBatchTransformDPPDetails {}; + template <> + struct DivergentBatchTransformDPPDetails {}; + template <> struct DivergentBatchTransformDPPDetails { uint numPlanes; }; -#if defined(__NVCC__) +#if defined(__NVCC__) || defined(__HIPCC__) template struct DivergentBatchTransformDPP { private: @@ -352,18 +358,33 @@ namespace fk { // namespace FusedKernel static constexpr ParArch PAR_ARCH = ParArch::GPU_NVIDIA; template FK_DEVICE_FUSE void exec(const DPPDetails&, const IOpSequenceTypes&... iOpSequences) { + const int x = (blockDim.x * blockIdx.x) + threadIdx.x; + const int y = (blockDim.y * blockIdx.y) + threadIdx.y; + const int z = blockIdx.z; + const Point thread{ x, y, z }; - const cg::thread_block g = cg::this_thread_block(); + Parent::template divergent_operate<0>(thread, iOpSequences...); + } + }; - const int x = (g.dim_threads().x * g.group_index().x) + g.thread_index().x; - const int y = (g.dim_threads().y * g.group_index().y) + g.thread_index().y; - const int z = g.group_index().z; + template + struct DivergentBatchTransformDPP { + private: + using Parent = DivergentBatchTransformDPPBase; + public: + using DPPDetails = DivergentBatchTransformDPPDetails; + static constexpr ParArch PAR_ARCH = ParArch::GPU_AMD; + template + FK_DEVICE_FUSE void exec(const DPPDetails&, const IOpSequenceTypes&... iOpSequences) { + const int x = (blockDim.x * blockIdx.x) + threadIdx.x; + const int y = (blockDim.y * blockIdx.y) + threadIdx.y; + const int z = blockIdx.z; const Point thread{ x, y, z }; Parent::template divergent_operate<0>(thread, iOpSequences...); } }; -#endif // defined(__NVCC__) +#endif // defined(__NVCC__) || defined(__HIPCC__) template struct DivergentBatchTransformDPP { private: diff --git a/include/fused_kernel/core/execution_model/executor_details/executor_kernels.h b/include/fused_kernel/core/execution_model/executor_details/executor_kernels.h index 41348be9..aee34599 100644 --- a/include/fused_kernel/core/execution_model/executor_details/executor_kernels.h +++ b/include/fused_kernel/core/execution_model/executor_details/executor_kernels.h @@ -15,20 +15,26 @@ #ifndef FK_EXECUTOR_KERNELS_H #define FK_EXECUTOR_KERNELS_H -#if defined(__NVCC__) +#if defined(__NVCC__) || defined(__HIPCC__) namespace fk { +#if defined(__NVCC__) +#define FK_GRID_CONSTANT __grid_constant__ +#else +#define FK_GRID_CONSTANT +#endif template -__global__ void launchDivergentBatchTransformDPP_Kernel(const __grid_constant__ DPPDetails details, - const __grid_constant__ IOpSequences... iOpSequences) { +__global__ void launchDivergentBatchTransformDPP_Kernel(const FK_GRID_CONSTANT DPPDetails details, + const FK_GRID_CONSTANT IOpSequences... iOpSequences) { DivergentBatchTransformDPP::exec(details, iOpSequences...); } template -__global__ void launchTransformDPP_Kernel(const __grid_constant__ TDPPDetails tDPPDetails, - const __grid_constant__ IOps... operations) { +__global__ void launchTransformDPP_Kernel(const FK_GRID_CONSTANT TDPPDetails tDPPDetails, + const FK_GRID_CONSTANT IOps... operations) { TransformDPP::exec(tDPPDetails, operations...); } +#undef FK_GRID_CONSTANT } // namespace fk #endif diff --git a/include/fused_kernel/core/execution_model/executors.h b/include/fused_kernel/core/execution_model/executors.h index d0f93839..0d11ee51 100644 --- a/include/fused_kernel/core/execution_model/executors.h +++ b/include/fused_kernel/core/execution_model/executors.h @@ -23,13 +23,13 @@ #include #include -#if defined(__NVCC__) +#if defined(__NVCC__) || defined(__HIPCC__) #include #endif namespace fk { -#if defined(__NVCC__) +#if defined(__NVCC__) || defined(__HIPCC__) struct CtxDim3 { uint x; uint y; @@ -147,16 +147,18 @@ FK_HOST_FUSE void executeOperations(const std::array, Batch>& input, co struct Executor { FK_STATIC_STRUCT(Executor, Executor) static_assert(DataParallelPattern::PAR_ARCH == ParArch::GPU_NVIDIA || + DataParallelPattern::PAR_ARCH == ParArch::GPU_AMD || DataParallelPattern::PAR_ARCH == ParArch::CPU || - DataParallelPattern::PAR_ARCH == ParArch::GPU_NVIDIA_JIT, "Only GPU_NVIDIA, CPU and GPU_NVIDIA_JIT are supported"); + DataParallelPattern::PAR_ARCH == ParArch::GPU_NVIDIA_JIT, "Only GPU_NVIDIA, GPU_AMD, CPU and GPU_NVIDIA_JIT are supported"); }; #else template struct Executor { FK_STATIC_STRUCT(Executor, Executor) static_assert(DataParallelPattern::PAR_ARCH == ParArch::GPU_NVIDIA || + DataParallelPattern::PAR_ARCH == ParArch::GPU_AMD || DataParallelPattern::PAR_ARCH == ParArch::CPU, - "Only GPU_NVIDIA and CPU supported"); + "Only GPU_NVIDIA, GPU_AMD and CPU supported"); }; #endif @@ -238,7 +240,7 @@ FK_HOST_FUSE void executeOperations(const std::array, Batch>& input, co } }; -#if defined(__NVCC__) +#if defined(__NVCC__) || defined(__HIPCC__) struct ComputeBestSolutionBase { FK_HOST_FUSE uint computeDiscardedThreads(const uint width, const uint height, const uint blockDimx, const uint blockDimy) { const uint modX = width % blockDimx; @@ -409,7 +411,110 @@ FK_HOST_FUSE void executeOperations(const std::array, Batch>& input, co executeOperations_helper(stream, iOpSequences...); } }; -#endif +#endif // __NVCC__ + +#if defined(__HIPCC__) + template + struct Executor> { + private: + using Child = Executor>; + using Parent = BaseExecutor; + template + FK_HOST_FUSE void executeOperations_helper(Stream_& stream_, const IOps&... iOps) { + const hipStream_t stream = stream_.getHIPStream(); + constexpr ParArch PA = ParArch::GPU_AMD; + const auto tDetails = TransformDPP::build_details(iOps...); + if constexpr (decltype(tDetails)::TFI::ENABLED) { + const ActiveThreads activeThreads = tDetails.activeThreads; + + const CtxDim3 ctx_block = getDefaultBlockSize(activeThreads.x, activeThreads.y); + + const dim3 block{ ctx_block.x, ctx_block.y, 1 }; + const dim3 grid{ static_cast(ceil(activeThreads.x / static_cast(block.x))), + static_cast(ceil(activeThreads.y / static_cast(block.y))), + activeThreads.z }; + if (!tDetails.threadDivisible) { + launchTransformDPP_Kernel<<>>(tDetails, iOps...); + gpuErrchk(hipGetLastError()); + } else { + launchTransformDPP_Kernel<<>>(tDetails, iOps...); + gpuErrchk(hipGetLastError()); + } + } else { + const auto readOp = get_arg<0>(iOps...); + + const ActiveThreads activeThreads = readOp.getActiveThreads(); + + const CtxDim3 ctx_block = getDefaultBlockSize(activeThreads.x, activeThreads.y); + + const dim3 block{ ctx_block.x, ctx_block.y, 1 }; + const dim3 grid{ static_cast(ceil(activeThreads.x / static_cast(block.x))), + static_cast(ceil(activeThreads.y / static_cast(block.y))), + activeThreads.z }; + launchTransformDPP_Kernel<<>>(tDetails, iOps...); + gpuErrchk(hipGetLastError()); + } + } + public: + FK_STATIC_STRUCT(Executor, Child) + FK_HOST_FUSE ParArch parArch() { + return ParArch::GPU_AMD; + } + DECLARE_EXECUTOR_PARENT_IMPL + }; + + template + struct Executor> { + private: + using DPPType = DivergentBatchTransformDPP; + using DPPDetails = typename DPPType::DPPDetails; + using SelfType = Executor; + + template + FK_HOST_FUSE ActiveThreads getActiveThreads(const IOpSequenceTypes&... iOpSequences) { + const uint x = cxp::max::f(get<0>(iOpSequences.iOps).getActiveThreads().x...); + const uint y = cxp::max::f(get<0>(iOpSequences.iOps).getActiveThreads().y...); + const uint z = cxp::sum::f(get<0>(iOpSequences.iOps).getActiveThreads().z...); + return ActiveThreads{ x, y, z }; + } + + template + FK_HOST_FUSE auto fuseBackSequence(const IOpSequence& iOpSeq) { + return buildOperationSequence_tup( + apply([](auto&&... args) { + return BackFuser::fuse_back(std::forward(args)...); + }, iOpSeq.iOps) + ); + } + + template + FK_HOST_FUSE void executeOperationsFused(Stream_& stream, const IOpSequenceTypes&... iOpSequences) { + const ActiveThreads activeThreads = getActiveThreads(iOpSequences...); + const DPPDetails details{}; + + const dim3 block(cxp::min::f(activeThreads.x, 32u), cxp::min::f(activeThreads.y, 8u)); + const dim3 grid(ceil(activeThreads.x / static_cast(block.x)), + ceil(activeThreads.y / static_cast(block.y)), activeThreads.z); + launchDivergentBatchTransformDPP_Kernel<<>>(details, iOpSequences...); + gpuErrchk(hipGetLastError()); + } + + template + FK_HOST_FUSE void executeOperations_helper(Stream_& stream, const IOpSequenceTypes&... iOpSequences) { + executeOperationsFused(stream, fuseBackSequence(iOpSequences)...); + } + + public: + FK_STATIC_STRUCT(Executor, SelfType) + FK_HOST_FUSE ParArch parArch() { + return ParArch::GPU_AMD; + } + template + FK_HOST_FUSE void executeOperations(Stream_& stream, const IOpSequenceTypes&... iOpSequences) { + executeOperations_helper(stream, iOpSequences...); + } + }; +#endif // __HIPCC__ } // namespace fk #endif // FK_EXECUTORS_CUH \ No newline at end of file diff --git a/include/fused_kernel/core/execution_model/parallel_architectures.h b/include/fused_kernel/core/execution_model/parallel_architectures.h index 487f6bf9..59553ad8 100644 --- a/include/fused_kernel/core/execution_model/parallel_architectures.h +++ b/include/fused_kernel/core/execution_model/parallel_architectures.h @@ -49,6 +49,8 @@ namespace fk { #if defined(__NVCC__) constexpr ParArch defaultParArch = ParArch::GPU_NVIDIA; +#elif defined(__HIPCC__) + constexpr ParArch defaultParArch = ParArch::GPU_AMD; #else constexpr ParArch defaultParArch = ParArch::CPU; #endif diff --git a/include/fused_kernel/core/execution_model/stream.h b/include/fused_kernel/core/execution_model/stream.h index 051670e1..e4ef285f 100644 --- a/include/fused_kernel/core/execution_model/stream.h +++ b/include/fused_kernel/core/execution_model/stream.h @@ -18,10 +18,8 @@ #include #include -#if defined(__NVCC__) +#if defined(__NVCC__) || defined(__HIPCC__) #include -#elif defined(__HIP__) -#include #endif namespace fk { @@ -109,6 +107,67 @@ namespace fk { }; #endif +#if defined(__HIPCC__) + template <> + class Stream_ final : public BaseStream { + hipStream_t m_stream; + bool m_isMine{ false }; + + inline void initFromOther(const Stream_& other) { + m_stream = other.m_stream; + m_isMine = other.m_isMine; + } + + public: + Stream_() : BaseStream() { + gpuErrchk(hipStreamCreate(&m_stream)); + m_isMine = true; + } + Stream_(const Stream_& other) : BaseStream(other) { + initFromOther(other); + } + explicit Stream_(const hipStream_t& stream) : m_stream(stream), BaseStream() {} + + hipStream_t operator()() const { + return m_stream; + } + + Stream_& operator=(const Stream_& other) { + if (this != &other) { + BaseStream::operator=(other); + initFromOther(other); + } + return *this; + } + + Stream_(Stream_&&) = delete; + Stream_& operator=(Stream_&&) = delete; + + ~Stream_() { + if (this->getRefCount() == 0 && m_stream != 0 && m_isMine) { + sync(); + gpuErrchk(hipStreamDestroy(m_stream)); + } + } + + operator hipStream_t() const { + return m_stream; + } + inline hipStream_t getHIPStream() const { + return m_stream; + } + inline void sync() final { + gpuErrchk(hipStreamSynchronize(m_stream)); + } + constexpr inline enum ParArch getParArch() const { + return ParArch::GPU_AMD; + }; + static constexpr inline enum ParArch parArch() { + return ParArch::GPU_AMD; + } + }; +#endif + template <> class Stream_ final : public BaseStream { public: diff --git a/include/fused_kernel/core/utils/compiler_macros.h.rej b/include/fused_kernel/core/utils/compiler_macros.h.rej deleted file mode 100644 index bf7354c7..00000000 --- a/include/fused_kernel/core/utils/compiler_macros.h.rej +++ /dev/null @@ -1,10 +0,0 @@ -diff a/include/fused_kernel/core/utils/compiler_macros.h b/include/fused_kernel/core/utils/compiler_macros.h (rejected hunks) -@@ -0,5 +0,7 @@ - #define CLANG_HOST_DEVICE 0 - #endif - -+#define VS2017_COMPILER (_MSC_VER_EXISTS && _MSC_VER >= 1910 && _MSC_VER < 1920) -+#define NO_VS2017_COMPILER !VS2017_COMPILER - - #endif // COMPILER_MACROS_H -\ No newline at end of file diff --git a/include/fused_kernel/core/utils/utils.h b/include/fused_kernel/core/utils/utils.h index 1629dc83..b31807b7 100644 --- a/include/fused_kernel/core/utils/utils.h +++ b/include/fused_kernel/core/utils/utils.h @@ -23,6 +23,8 @@ #if defined(__NVCC__) #include +#elif defined(__HIPCC__) +#include #endif #if defined(NVRTC_ENABLED) @@ -31,7 +33,7 @@ #endif #endif // NVRTC_COMPILER -#if defined(__NVCC__) +#if defined(__NVCC__) || defined(__HIPCC__) #define FK_DEVICE_FUSE __device__ __forceinline__ static constexpr #define FK_DEVICE_STATIC __device__ __forceinline__ static #define FK_DEVICE_CNST __device__ __forceinline__ constexpr @@ -95,8 +97,9 @@ using ulonglong = unsigned long long; using ushort = unsigned short; using ulong = unsigned long; -#if defined(__NVCC__) +#if defined(__NVCC__) || defined(__HIPCC__) namespace fk { +#if defined(__NVCC__) inline void gpuAssert(cudaError_t code, const char *file, int line, @@ -112,6 +115,23 @@ namespace fk { if (abort) throw std::runtime_error(message.c_str()); } } +#else + inline void gpuAssert(hipError_t code, + const char *file, + int line, + bool abort = true) { + if (code != hipSuccess) { + std::string message = "GPU Error: "; + message.append(hipGetErrorString(code)); + message.append(" File: "); + message.append(file); + message.append(" Line:"); + message.append(std::to_string(line).c_str()); + message.append("\n"); + if (abort) throw std::runtime_error(message.c_str()); + } + } +#endif // __NVCC__ #if defined(NVRTC_ENABLED) inline void gpuAssert(CUresult code, const char* file, @@ -150,7 +170,7 @@ namespace fk { } // namespace fk #define gpuErrchk(ans) { fk::gpuAssert((ans), __FILE__, __LINE__, true); } -#endif // defined(__NVCC__) +#endif // defined(__NVCC__) || defined(__HIPCC__) // Null type, used for Operation required aliases that can not still be known, // because they are deduced from a backwards operation that is till not defined. diff --git a/include/fused_kernel/core/utils/vector_utils.h b/include/fused_kernel/core/utils/vector_utils.h index 6dbe75a8..70188816 100644 --- a/include/fused_kernel/core/utils/vector_utils.h +++ b/include/fused_kernel/core/utils/vector_utils.h @@ -268,7 +268,8 @@ namespace fk { template FK_HOST_DEVICE_CNST T make_(const Numbers&... pack) { - if constexpr (std::is_aggregate_v) { + // Not is_aggregate_v: HIP vector types are non-aggregate class templates. + if constexpr (vector_type) { return make::type(pack...); } else { static_assert(sizeof...(pack) == 1, "passing more than one argument for a non cuda vector type"); @@ -278,14 +279,15 @@ namespace fk { template FK_HOST_DEVICE_CNST T make_set(const VBase val) { + // Direct-list-init via make::type: HIP's 1-arg vector ctor is explicit. if constexpr (cn == 1) { - return {val}; + return make::type(val); } else if constexpr (cn == 2) { - return {val, val}; + return make::type(val, val); } else if constexpr (cn == 3) { - return {val, val, val}; + return make::type(val, val, val); } else { - return {val, val, val, val}; + return make::type(val, val, val, val); } } diff --git a/lib/CMakeLists.txt b/lib/CMakeLists.txt index 393e4a38..901116e0 100644 --- a/lib/CMakeLists.txt +++ b/lib/CMakeLists.txt @@ -3,7 +3,7 @@ add_library(${LIB_NAME} INTERFACE) # add alias so the project can be uses with add_subdirectory add_library(${LIB_NAME}::${LIB_NAME} ALIAS ${LIB_NAME}) - target_include_directories( +target_include_directories( ${LIB_NAME} INTERFACE $ $) diff --git a/tests/data/basic_test.h b/tests/data/basic_test.h index 85521ddd..beaa4a91 100644 --- a/tests/data/basic_test.h +++ b/tests/data/basic_test.h @@ -123,10 +123,10 @@ int launch() { std::is_same_v, fk::Write>>; if (test2Dpassed && fk::and_v) { - std::cout << "cuda_transform executed!!" << std::endl; + std::cout << "gpu_transform executed!!" << std::endl; return 0; } else { - std::cout << "cuda_transform failed!!" << std::endl; + std::cout << "gpu_transform failed!!" << std::endl; if (!test2Dpassed) { std::cout << "Specifically testPtr_2D failed!!" << std::endl; } diff --git a/tests/data/test_ptr_nd.h b/tests/data/test_ptr_nd.h index a07d9427..f0fa2724 100644 --- a/tests/data/test_ptr_nd.h +++ b/tests/data/test_ptr_nd.h @@ -149,7 +149,7 @@ int launch() { bool h_correct{ true }; for (int y = 0; y < HEIGHT; y++) { for (int x = 0; x < WIDTH; x++) { - const Bool3 boolVect = *PtrAccessor::cr_point(Point{x, y, 0}, test0.ptrPinned()) == make_(1, 2, 3); + const bool boolVect = *PtrAccessor::cr_point(Point{x, y, 0}, test0.ptrPinned()) == make_(1, 2, 3); const bool allTrue = boolVect; h_correct &= allTrue; } diff --git a/tests/launcher.in b/tests/launcher.in index 6ce39533..27da003d 100644 --- a/tests/launcher.in +++ b/tests/launcher.in @@ -1,5 +1,5 @@ /* Copyright 2023 Oscar Amoros Huguet - Copyright 2023 Albert Andaluz Gonzalez + Copyright 2023-2026 Albert Andaluz Gonzalez Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except in compliance with the License. diff --git a/utests/algorithm/image_processing/utest_color_conversion.h b/utests/algorithm/image_processing/utest_color_conversion.h index 00267ecc..ec728100 100644 --- a/utests/algorithm/image_processing/utest_color_conversion.h +++ b/utests/algorithm/image_processing/utest_color_conversion.h @@ -289,7 +289,7 @@ void testSaturateDenormalizePixel() { }; std::array expectedVals = { - uchar3{0, 128, 255}, // Saturated to [0,1] then denormalized to [0,255] + uchar3{0, 127, 255}, // Saturated to [0,1] then denormalized to [0,255]; 0.5*255=127.5 truncates to 127 uchar3{63, 191, 229} // 0.25*255=63.75≈63, 0.75*255=191.25≈191, 0.9*255=229.5≈229 }; @@ -553,11 +553,19 @@ constexpr M3x3Float // ========================================================================= // 10-BIT and 12-BIT LIMITED RANGE matrices // ========================================================================= +template +constexpr M3x3Float scaleLimitedMatrix(const M3x3Float& matrix) { + const M3x3Float& factors = limitedFactors; + return {{matrix.x.x * factors.x.x, matrix.x.y * factors.x.y, matrix.x.z * factors.x.z}, + {matrix.y.x * factors.y.x, matrix.y.y * factors.y.y, matrix.y.z * factors.y.z}, + {matrix.z.x * factors.z.x, matrix.z.y * factors.z.y, matrix.z.z * factors.z.z}}; +} + template requires one_of_v, TypeList, CD_t>> constexpr M3x3Float ccMatrixTest = - ccMatrixTest * limitedFactors; + scaleLimitedMatrix(ccMatrixTest); // Normalized float (fn) matrices share the exact same scaling ratio as their integer equivalents. template @@ -570,7 +578,11 @@ template constexpr M3x3Float ccMatrixTest = ccMatrixTest; constexpr bool equalM3x3(const M3x3Float& first, const M3x3Float& second) { - return (first.x == second.x) && (first.y == second.y) && (first.z == second.z); + // Avoid the vector equality operator here: it is not constexpr on all + // supported compilers, while the scalar comparisons are. + return first.x.x == second.x.x && first.x.y == second.x.y && first.x.z == second.x.z && + first.y.x == second.y.x && first.y.y == second.y.y && first.y.z == second.y.z && + first.z.x == second.z.x && first.z.y == second.z.y && first.z.z == second.z.z; } template diff --git a/utests/algorithm/image_processing/utest_saturate/CMakeLists.txt b/utests/algorithm/image_processing/utest_saturate/CMakeLists.txt index ea9fba20..e7b5e6d1 100644 --- a/utests/algorithm/image_processing/utest_saturate/CMakeLists.txt +++ b/utests/algorithm/image_processing/utest_saturate/CMakeLists.txt @@ -44,6 +44,29 @@ if (CMAKE_CUDA_COMPILER AND ENABLE_CUDA) endforeach() endif() + +if (CMAKE_HIP_COMPILER AND ENABLE_HIP) + set(EXTENSION hip) + string(TOUPPER ${EXTENSION} EXTENSION_UPPER) + configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${TARGET_NAME}.h.in + ${CMAKE_BINARY_DIR}/${TARGET_NAME}/${TARGET_NAME}_${EXTENSION}.h) + + add_generated_test_stub("${TARGET_NAME}_${EXTENSION}" + "${CMAKE_BINARY_DIR}/${TARGET_NAME}/${TARGET_NAME}_${EXTENSION}.h" + "") + set_ide_target_folder(${TARGET_NAME}_${EXTENSION} "/utests/algorithm/image_processing" "${EXTENSION}") + + add_hip_to_test("${TARGET_NAME}_${EXTENSION}") + foreach(FUNDAMENTAL_TYPE ${FUNDAMENTAL_TYPES}) + add_shared_test_lib(${TARGET_NAME} + "" + "${EXTENSION}" + "${FUNDAMENTAL_TYPE}" + ) + set_property(TARGET "${TARGET_NAME}_${FUNDAMENTAL_TYPE}_${EXTENSION}" PROPERTY FOLDER "/utests/algorithm/image_processing/hip/${TARGET_NAME}") + target_link_libraries(${TARGET_NAME}_${EXTENSION} PRIVATE "${TARGET_NAME}_${FUNDAMENTAL_TYPE}_${EXTENSION}") + endforeach() +endif() \ No newline at end of file diff --git a/utests/algorithm/image_processing/utest_saturate/utest_saturate_ftype.hip.in b/utests/algorithm/image_processing/utest_saturate/utest_saturate_ftype.hip.in new file mode 100644 index 00000000..5b948a34 --- /dev/null +++ b/utests/algorithm/image_processing/utest_saturate/utest_saturate_ftype.hip.in @@ -0,0 +1,26 @@ +/* Copyright 2025-2026 Oscar Amoros Huguet + Copyright 2025-2026 Grup Mediapro S.L.U + Copyright 2026 Albert Andaluz González + Licensed under the Apache License, Version 2.0 (the "License"); + you may not use this file except in compliance with the License. + You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + + Unless required by applicable law or agreed to in writing, software + distributed under the License is distributed on an "AS IS" BASIS, + WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + See the License for the specific language governing permissions and + limitations under the License. */ + +#include "utest_saturate_@FUNDAMENTAL_TYPE@_@EXTENSION@.h" + +START_ADDING_TESTS +using Fundamental = fk::RemoveType_t<0, fk::StandardTypes>; +addAllOutputTestsForInput(std::make_index_sequence{}); +STOP_ADDING_TESTS + +int fk::utest_saturate_@FUNDAMENTAL_TYPE@_@EXTENSION@::launch() { + RUN_ALL_TESTS + return 0; +}