From adece8c6ed74c248955eec5b78079a01840bfa73 Mon Sep 17 00:00:00 2001 From: Viktor Yastrebov Date: Mon, 17 Jun 2024 17:41:32 +0300 Subject: [PATCH] reimplement external intrinsic idea to make it embeded into final transpiled source code --- lib/core/intrinsics/external_intrinsics.cpp | 102 ++++++------------ lib/core/intrinsics/external_intrinsics.h | 20 +--- .../transpiler_session/code_generator.cpp | 22 ++-- lib/core/transpiler_session/header_info.cpp | 83 ++++---------- lib/core/transpiler_session/header_info.h | 19 +--- lib/core/transpiler_session/session_stage.cpp | 3 + .../stages/transpiler/transpilation.cpp | 2 +- .../data/misc/custom_intrinsics/api.h | 4 + .../cuda/{impl.h => impl.cpp} | 4 +- .../dpcpp/{impl.h => impl.cpp} | 9 +- .../data/misc/custom_intrinsics/hip/impl.cpp | 7 ++ .../data/misc/custom_intrinsics/hip/impl.h | 8 -- .../launcher/{impl.h => impl.cpp} | 2 - .../{serial/impl.h => openmp/impl.cpp} | 4 +- .../{openmp/impl.h => serial/impl.cpp} | 2 - .../user_intrinsics/user_intrinsics_ref.cpp | 1 + .../user_intrinsics/user_intrinsics_ref.cpp | 3 +- .../user_intrinsics/user_intrinsics_ref.cpp | 3 +- .../user_intrinsics/user_intrinsics_ref.cpp | 6 +- .../user_intrinsics/user_intrinsics_ref.cpp | 6 +- 20 files changed, 103 insertions(+), 207 deletions(-) rename tests/functional/data/misc/custom_intrinsics/cuda/{impl.h => impl.cpp} (68%) rename tests/functional/data/misc/custom_intrinsics/dpcpp/{impl.h => impl.cpp} (57%) create mode 100644 tests/functional/data/misc/custom_intrinsics/hip/impl.cpp delete mode 100644 tests/functional/data/misc/custom_intrinsics/hip/impl.h rename tests/functional/data/misc/custom_intrinsics/launcher/{impl.h => impl.cpp} (90%) rename tests/functional/data/misc/custom_intrinsics/{serial/impl.h => openmp/impl.cpp} (64%) rename tests/functional/data/misc/custom_intrinsics/{openmp/impl.h => serial/impl.cpp} (88%) diff --git a/lib/core/intrinsics/external_intrinsics.cpp b/lib/core/intrinsics/external_intrinsics.cpp index b54fd3d9..46d2770a 100644 --- a/lib/core/intrinsics/external_intrinsics.cpp +++ b/lib/core/intrinsics/external_intrinsics.cpp @@ -3,6 +3,7 @@ #include #include "core/transpiler_session/session_stage.h" #include "core/transpiler_session/transpiler_session.h" +#include "oklt/util/io_helper.h" #include "util/string_utils.hpp" #include @@ -73,10 +74,9 @@ std::optional getExternalInstrincisInclude(TranspilerSession& session, return std::nullopt; } -tl::expected, std::string> getExternalIntrinsicSource( - TargetBackend backend, - const fs::path& intrinsicPath, - clang::SourceManager& sm) { +tl::expected getExternalIntrinsicSource(TargetBackend backend, + const fs::path& intrinsicPath, + clang::SourceManager& sm) { auto implPathResult = getIntrincisImplSourcePath(backend, intrinsicPath); if (!implPathResult) { return tl::make_unexpected(implPathResult.error()); @@ -93,7 +93,7 @@ tl::expected, std::string> getExternalIntrin } auto it = std::find_if(files.cbegin(), files.cend(), [](const fs::path& p) -> bool { - return p.extension().string() == std::string(".h"); + return p.extension().string() == std::string(".cpp"); }); if (it == files.cend()) { @@ -101,25 +101,20 @@ tl::expected, std::string> getExternalIntrin return tl::make_unexpected(error); } - auto& fm = sm.getFileManager(); - auto backendFilePath = it->string(); - auto maybeReplacedFile = fm.getFile(backendFilePath); - if (!maybeReplacedFile) { - std::string error = "Can't open file: " + backendFilePath; + auto contentResult = util::readFileAsStr(*it); + if (!contentResult) { + std::string error = "Can't get memory buffer for: " + it->string(); return tl::make_unexpected(error); } - auto maybeBuffer = fm.getBufferForFile(maybeReplacedFile.get()); - if (!maybeBuffer) { - std::string error = "Can't get memory buffer for: " + backendFilePath; - return tl::make_unexpected(error); - } - return std::move(maybeBuffer.get()); + return contentResult.value(); } -bool overrideExternalIntrinsic(TranspilerSession& session, +bool overrideExternalIntrinsic(SessionStage& stage, + HeaderDepsInfo& deps, const std::string& includedFileName, - clang::OptionalFileEntryRef includedFile, - clang::SourceManager& sourceManager) { + clang::OptionalFileEntryRef includedFile) { + auto& session = stage.getSession(); + auto& sourceManager = stage.getCompiler().getSourceManager(); const auto& userIntrinsics = session.getInput().userIntrinsics; if (!userIntrinsics.empty()) { auto maybeIntrinsicPath = getExternalInstrincisInclude(session, includedFileName); @@ -127,25 +122,27 @@ bool overrideExternalIntrinsic(TranspilerSession& session, return false; } auto intrinsicPath = maybeIntrinsicPath.value(); - auto infoResult = - getExternalIntrinsicSource(session.getInput().backend, intrinsicPath, sourceManager); - if (!infoResult) { - session.pushError(std::error_code(), infoResult.error()); + auto intrinsicResult = + getExternalIntrinsicSource(stage.getBackend(), intrinsicPath, sourceManager); + if (!intrinsicResult) { + session.pushError(std::error_code(), intrinsicResult.error()); return false; } - auto buffer = std::move(infoResult.value()); - auto& fm = sourceManager.getFileManager(); + deps.externalIntrinsicsSources[includedFileName] = std::move(intrinsicResult.value()); + + auto emptyExternalIntrinsic = MemoryBuffer::getMemBuffer(""); if (includedFile) { auto fileRef = includedFile; const auto& fileEntry = fileRef->getFileEntry(); - sourceManager.overrideFileContents(&fileEntry, std::move(buffer)); + sourceManager.overrideFileContents(&fileEntry, std::move(emptyExternalIntrinsic)); } else { // INFO: case when the file can be found by relative path // it happens when the include path is relative to WORKING DIR path + auto& fm = sourceManager.getFileManager(); auto maybeFileRef = fm.getFileRef(includedFileName); if (maybeFileRef) { auto foundFileRef = maybeFileRef.get(); - sourceManager.overrideFileContents(foundFileRef, std::move(buffer)); + sourceManager.overrideFileContents(foundFileRef, std::move(emptyExternalIntrinsic)); } } return true; @@ -153,59 +150,28 @@ bool overrideExternalIntrinsic(TranspilerSession& session, return false; } -void nullyLauncherExternalIntrinsics(TransformedFiles& inputs, SessionStage& stage) { - if (stage.getBackend() != TargetBackend::_LAUNCHER) { - return; - } - - auto& session = stage.getSession(); - const auto& intrinsics = session.getInput().userIntrinsics; - if (intrinsics.empty()) { +void updateExternalIntrinsicMap(SessionStage& stage, HeaderDepsInfo& deps) { + if (deps.externalIntrinsicsSources.empty()) { return; } - for (auto& mappedFile : inputs.fileMap) { - auto fileName = normalizedFileName(mappedFile.first); - for (const auto& includePath : intrinsics) { - auto folderPrefix = includePath.filename().string(); - if (util::startsWith(fileName, folderPrefix)) { - mappedFile.second.clear(); - } - } - } -} - -void embedLauncherExternalIntrinsics(std::string& input, - const HeaderDepsInfo& info, - SessionStage& stage) { auto backend = stage.getBackend(); - if (backend != TargetBackend::_LAUNCHER) { - return; - } - - const auto& usedIntrinsics = info.externalIntrinsics; - if (usedIntrinsics.empty()) { - return; - } - - auto session = stage.getSession(); + auto& session = stage.getSession(); auto& sm = stage.getCompiler().getSourceManager(); - for (const auto& includeIntrinsic : usedIntrinsics) { - auto maybeIntrinsicPath = getExternalInstrincisInclude(session, includeIntrinsic); + for (auto& mappedIntrinsic : deps.externalIntrinsicsSources) { + auto maybeIntrinsicPath = getExternalInstrincisInclude(session, mappedIntrinsic.first); if (!maybeIntrinsicPath) { - std::string error = "Count not find implementation for " + includeIntrinsic; + std::string error = "Count not find implementation for " + mappedIntrinsic.first; session.pushError(std::error_code(), error); return; } auto intrinsicPath = maybeIntrinsicPath.value(); - auto infoResult = getExternalIntrinsicSource(backend, intrinsicPath, sm); - if (!infoResult) { - session.pushError(std::error_code(), infoResult.error()); + auto intrinsicResult = getExternalIntrinsicSource(backend, intrinsicPath, sm); + if (!intrinsicResult) { + session.pushError(std::error_code(), intrinsicResult.error()); return; } - auto buffer = std::move(infoResult.value()); - // INFO: make temporary buffer because pointer could be not null terminated - input.insert(0, buffer->getBuffer().str()); + mappedIntrinsic.second = std::move(intrinsicResult.value()); } } } // namespace oklt diff --git a/lib/core/intrinsics/external_intrinsics.h b/lib/core/intrinsics/external_intrinsics.h index 4b852a11..4ad4a8bd 100644 --- a/lib/core/intrinsics/external_intrinsics.h +++ b/lib/core/intrinsics/external_intrinsics.h @@ -2,27 +2,15 @@ #include #include -namespace clang { -class CompilerInstance; -class SourceManager; -} // namespace clang - namespace oklt { -class TranspilerSession; -class TransformedFiles; class SessionStage; class HeaderDepsInfo; -bool overrideExternalIntrinsic(TranspilerSession& session, +bool overrideExternalIntrinsic(SessionStage& stage, + HeaderDepsInfo& deps, const std::string& includedFileName, - clang::OptionalFileEntryRef includedFile, - clang::SourceManager& sourceManager); - -void nullyLauncherExternalIntrinsics(TransformedFiles& inputs, SessionStage& stage); - -void embedLauncherExternalIntrinsics(std::string& input, - const HeaderDepsInfo& info, - SessionStage& stage); + clang::OptionalFileEntryRef includedFile); +void updateExternalIntrinsicMap(SessionStage& stage, HeaderDepsInfo& deps); } // namespace oklt diff --git a/lib/core/transpiler_session/code_generator.cpp b/lib/core/transpiler_session/code_generator.cpp index a0539585..2b905ed4 100644 --- a/lib/core/transpiler_session/code_generator.cpp +++ b/lib/core/transpiler_session/code_generator.cpp @@ -7,11 +7,11 @@ #include "core/transpiler_session/transpiler_session.h" #include "core/intrinsics/builtin_intrinsics.h" -#include "core/intrinsics/external_intrinsics.h" +// #include "core/intrinsics/external_intrinsics.h" #include "core/handler_manager/handler_manager.h" -#include "core/utils/attributes.h" +// #include "core/utils/attributes.h" #include "core/vfs/overlay_fs.h" #include @@ -86,10 +86,8 @@ void removeSystemHeaders(SessionStage& stage, const HeaderDepsInfo& deps) { rewriter.RemoveText({dep.hashLoc, dep.filenameRange.getEnd()}); } - for (const auto& intrinsicDep : deps.externalIntrinsicDeps) { - SPDLOG_TRACE( - "remove system include {} {}", intrinsicDep.relativePath, intrinsicDep.fileName); - rewriter.RemoveText({intrinsicDep.hashLoc, intrinsicDep.filenameRange.getEnd()}); + for (const auto& intrinsic : deps.externalIntrinsicHeaders) { + rewriter.RemoveText({intrinsic.hashLoc, intrinsic.filenameRange.getEnd()}); } } @@ -169,18 +167,14 @@ std::string restoreSystemAndBackendHeaders(SessionStage& stage, } } - embedLauncherExternalIntrinsics(input, deps, stage); + for (const auto& externalIntrinsicSource : deps.externalIntrinsicsSources) { + input.insert(0, externalIntrinsicSource.second); + } for (auto it = deps.backendHeaders.rbegin(); it < deps.backendHeaders.rend(); ++it) { input.insert(0, *it); } - if (backend != TargetBackend::_LAUNCHER) { - for (const auto& intrinsicDep : deps.externalIntrinsicDeps) { - input.insert(0, "#include <" + intrinsicDep.fileName + ">\n"); - } - } - // restore system headers for (const auto& dep : deps.topLevelDeps) { if (!clang::SrcMgr::isSystem(dep.fileType)) { @@ -197,8 +191,6 @@ tl::expected fuseIncludeDeps(SessionStage& stage, const Head auto inputs = gatherTransformedFiles(stage); - nullyLauncherExternalIntrinsics(inputs, stage); - auto preprocessedResult = preprocesseInputs(stage, inputs); if (!preprocessedResult) { return preprocessedResult; diff --git a/lib/core/transpiler_session/header_info.cpp b/lib/core/transpiler_session/header_info.cpp index 33a9e6a6..a1d95d90 100644 --- a/lib/core/transpiler_session/header_info.cpp +++ b/lib/core/transpiler_session/header_info.cpp @@ -1,20 +1,19 @@ #include "core/transpiler_session/header_info.h" #include "core/intrinsics/builtin_intrinsics.h" #include "core/intrinsics/external_intrinsics.h" -#include "core/transpiler_session/transpiler_session.h" +#include "core/transpiler_session/session_stage.h" namespace {} namespace oklt { using namespace llvm; -InclusionDirectiveCallback::InclusionDirectiveCallback(TranspilerSession& session, +InclusionDirectiveCallback::InclusionDirectiveCallback(SessionStage& stage, HeaderDepsInfo& deps_, clang::SourceManager& sm_) : deps(deps_), sm(sm_), - _session(session), - _isInExternalIntrinsic(false) {} + _stage(stage) {} void InclusionDirectiveCallback::InclusionDirective(clang::SourceLocation hashLoc, const clang::Token& includeTok, @@ -36,67 +35,25 @@ void InclusionDirectiveCallback::InclusionDirective(clang::SourceLocation hashLo } auto fileNameStr = fileName.str(); - auto isIntrinsic = overrideExternalIntrinsic(_session, fileNameStr, file, sm); - if (isIntrinsic) { - deps.externalIntrinsics.push_back(fileNameStr); - _extIntrinsicFID = sm.getFileID(hashLoc); - } + bool isIntrinsic = overrideExternalIntrinsic(_stage, deps, fileNameStr, file); + + auto dep = HeaderDep{ + .hashLoc = hashLoc, + .includeTok = includeTok, + .fileName = fileNameStr, + .isAngled = isAngled, + .filenameRange = filenameRange, + .file = file, + .searchPath = searchPath.str(), + .relativePath = relativePath.str(), + .imported = imported, + .fileType = fileType, + }; - // INFO: force instrinsic includes to be system includes - // they will be removed & attached to final transpiled source - if (_isInExternalIntrinsic) { - fileType = clang::SrcMgr::CharacteristicKind::C_System; - deps.externalIntrinsicDeps.push_back(HeaderDep{ - .hashLoc = hashLoc, - .includeTok = includeTok, - .fileName = fileNameStr, - .isAngled = isAngled, - .filenameRange = filenameRange, - .file = file, - .searchPath = searchPath.str(), - .relativePath = relativePath.str(), - .imported = imported, - .fileType = fileType, - }); + if (isIntrinsic) { + deps.externalIntrinsicHeaders.push_back(std::move(dep)); } else { - deps.topLevelDeps.push_back(HeaderDep{ - .hashLoc = hashLoc, - .includeTok = includeTok, - .fileName = fileNameStr, - .isAngled = isAngled, - .filenameRange = filenameRange, - .file = file, - .searchPath = searchPath.str(), - .relativePath = relativePath.str(), - .imported = imported, - .fileType = fileType, - }); + deps.topLevelDeps.push_back(std::move(dep)); } } - -void InclusionDirectiveCallback::FileChanged(clang::SourceLocation Loc, - FileChangeReason Reason, - clang::SrcMgr::CharacteristicKind FileType, - clang::FileID PrevFID) { - if (Reason == FileChangeReason::EnterFile) { - if (_extIntrinsicFID.isValid()) { - _isInExternalIntrinsic = PrevFID == _extIntrinsicFID; - return; - } - } - - if (Reason == FileChangeReason::ExitFile) { - auto thisFID = sm.getFileID(Loc); - if (_extIntrinsicFID.isValid() && thisFID == _extIntrinsicFID) { - _extIntrinsicFID = clang::FileID(); - _isInExternalIntrinsic = false; - return; - } - } -} - -bool InclusionDirectiveCallback::FileNotFound(clang::StringRef FileName) { - return _isInExternalIntrinsic; -} - } // namespace oklt diff --git a/lib/core/transpiler_session/header_info.h b/lib/core/transpiler_session/header_info.h index e64651d9..815334fd 100644 --- a/lib/core/transpiler_session/header_info.h +++ b/lib/core/transpiler_session/header_info.h @@ -30,16 +30,16 @@ struct HeaderDepsInfo { std::vector topLevelDeps; std::vector backendHeaders; std::vector backendNss; - std::vector externalIntrinsics; - std::vector externalIntrinsicDeps; + std::vector externalIntrinsicHeaders; + std::map externalIntrinsicsSources; bool useOklIntrinsic = false; }; -class TranspilerSession; +class SessionStage; class InclusionDirectiveCallback : public clang::PPCallbacks { public: - InclusionDirectiveCallback(TranspilerSession& session, + InclusionDirectiveCallback(SessionStage& session, HeaderDepsInfo& depsInfo, clang::SourceManager& sm); void InclusionDirective(clang::SourceLocation HashLoc, @@ -53,19 +53,10 @@ class InclusionDirectiveCallback : public clang::PPCallbacks { const clang::Module* Imported, clang::SrcMgr::CharacteristicKind FileType) override; - void FileChanged(clang::SourceLocation Loc, - FileChangeReason Reason, - clang::SrcMgr::CharacteristicKind FileType, - clang::FileID PrevFID = clang::FileID()) override; - - bool FileNotFound(clang::StringRef FileName) override; - private: HeaderDepsInfo& deps; clang::SourceManager& sm; - TranspilerSession& _session; - bool _isInExternalIntrinsic; - clang::FileID _extIntrinsicFID; + SessionStage& _stage; }; } // namespace oklt diff --git a/lib/core/transpiler_session/session_stage.cpp b/lib/core/transpiler_session/session_stage.cpp index 89bac2f7..89244531 100644 --- a/lib/core/transpiler_session/session_stage.cpp +++ b/lib/core/transpiler_session/session_stage.cpp @@ -1,6 +1,7 @@ #include "core/transpiler_session/session_stage.h" #include "core/diag/diag_consumer.h" #include "core/handler_manager/handler_manager.h" +#include "core/intrinsics/external_intrinsics.h" #include "core/transpiler_session/transpiler_session.h" #include @@ -33,6 +34,8 @@ void SessionStage::setLauncherMode() { _rewriter = std::make_unique(_compiler.getSourceManager(), _compiler.getLangOpts()); _backend = TargetBackend::_LAUNCHER; + auto& deps = tryEmplaceUserCtx(); + updateExternalIntrinsicMap(*this, deps); } std::string SessionStage::getRewriterResultForMainFile() { diff --git a/lib/pipeline/stages/transpiler/transpilation.cpp b/lib/pipeline/stages/transpiler/transpilation.cpp index 4d1aaf1a..744ecaff 100644 --- a/lib/pipeline/stages/transpiler/transpilation.cpp +++ b/lib/pipeline/stages/transpiler/transpilation.cpp @@ -441,7 +441,7 @@ class Transpilation : public StageAction { auto& deps = _stage->tryEmplaceUserCtx(); std::unique_ptr callback = std::make_unique( - *_session, deps, compiler.getSourceManager()); + *_stage, deps, compiler.getSourceManager()); // setup preprocessor hook to gather all user/system includes compiler.getPreprocessor().addPPCallbacks(std::move(callback)); diff --git a/tests/functional/data/misc/custom_intrinsics/api.h b/tests/functional/data/misc/custom_intrinsics/api.h index e69de29b..043542e1 100644 --- a/tests/functional/data/misc/custom_intrinsics/api.h +++ b/tests/functional/data/misc/custom_intrinsics/api.h @@ -0,0 +1,4 @@ +#pragma once + +bool okl_is_nan(float value); + diff --git a/tests/functional/data/misc/custom_intrinsics/cuda/impl.h b/tests/functional/data/misc/custom_intrinsics/cuda/impl.cpp similarity index 68% rename from tests/functional/data/misc/custom_intrinsics/cuda/impl.h rename to tests/functional/data/misc/custom_intrinsics/cuda/impl.cpp index f1afd79b..3e13fdb8 100644 --- a/tests/functional/data/misc/custom_intrinsics/cuda/impl.h +++ b/tests/functional/data/misc/custom_intrinsics/cuda/impl.cpp @@ -1,7 +1,5 @@ -#pragma once - //INFO: transpiling to cuda backend already include necessary header -bool okl_is_nan(float value) { +__device__ bool okl_is_nan(float value) { return isnan(value) != 0; } diff --git a/tests/functional/data/misc/custom_intrinsics/dpcpp/impl.h b/tests/functional/data/misc/custom_intrinsics/dpcpp/impl.cpp similarity index 57% rename from tests/functional/data/misc/custom_intrinsics/dpcpp/impl.h rename to tests/functional/data/misc/custom_intrinsics/dpcpp/impl.cpp index 9b698d21..ae44d84c 100644 --- a/tests/functional/data/misc/custom_intrinsics/dpcpp/impl.h +++ b/tests/functional/data/misc/custom_intrinsics/dpcpp/impl.cpp @@ -1,11 +1,8 @@ -#pragma once - -//INFO: from documentation +// INFO: from documentation // isNaN // Description:The function returns 1, if and only if its argument is a NaN. // Calling interface: -// int __binary32_isNaN(float x); - -bool okl_is_nan(float value) { +// int __binary32_isNaN(float x); +SYCL_EXTERNAL bool okl_is_nan(float value) { return __binary32_isNaN(value) == 1; } diff --git a/tests/functional/data/misc/custom_intrinsics/hip/impl.cpp b/tests/functional/data/misc/custom_intrinsics/hip/impl.cpp new file mode 100644 index 00000000..19491ff8 --- /dev/null +++ b/tests/functional/data/misc/custom_intrinsics/hip/impl.cpp @@ -0,0 +1,7 @@ +// INFO: relate to the documentation is must be supported natively +// https://rocm.docs.amd.com/projects/HIP/en/latest/reference/kernel_language.html#single-precision-mathematical-functions + +__device__ bool okl_is_nan(float value) { + return isnan(value); +} + diff --git a/tests/functional/data/misc/custom_intrinsics/hip/impl.h b/tests/functional/data/misc/custom_intrinsics/hip/impl.h deleted file mode 100644 index 36e4097e..00000000 --- a/tests/functional/data/misc/custom_intrinsics/hip/impl.h +++ /dev/null @@ -1,8 +0,0 @@ -#pragma once - -//INFO: relate to the documentation is must be supported natively -// https://rocm.docs.amd.com/projects/HIP/en/latest/reference/kernel_language.html#single-precision-mathematical-functions - -bool okl_is_nan(float value) { - return isnan(value); -} diff --git a/tests/functional/data/misc/custom_intrinsics/launcher/impl.h b/tests/functional/data/misc/custom_intrinsics/launcher/impl.cpp similarity index 90% rename from tests/functional/data/misc/custom_intrinsics/launcher/impl.h rename to tests/functional/data/misc/custom_intrinsics/launcher/impl.cpp index 8f43ea0a..42a0f9d0 100644 --- a/tests/functional/data/misc/custom_intrinsics/launcher/impl.h +++ b/tests/functional/data/misc/custom_intrinsics/launcher/impl.cpp @@ -1,5 +1,3 @@ -#pragma once - //INFO: needs to keep the old OCCA compatibility // some types must be knonw for launcher diff --git a/tests/functional/data/misc/custom_intrinsics/serial/impl.h b/tests/functional/data/misc/custom_intrinsics/openmp/impl.cpp similarity index 64% rename from tests/functional/data/misc/custom_intrinsics/serial/impl.h rename to tests/functional/data/misc/custom_intrinsics/openmp/impl.cpp index 9b7c6b51..5c88a0a1 100644 --- a/tests/functional/data/misc/custom_intrinsics/serial/impl.h +++ b/tests/functional/data/misc/custom_intrinsics/openmp/impl.cpp @@ -1,6 +1,4 @@ -#pragma once - -//INFO: needed for std::isnan +// INFO: needed for std::isnan #include bool okl_is_nan(float value) { diff --git a/tests/functional/data/misc/custom_intrinsics/openmp/impl.h b/tests/functional/data/misc/custom_intrinsics/serial/impl.cpp similarity index 88% rename from tests/functional/data/misc/custom_intrinsics/openmp/impl.h rename to tests/functional/data/misc/custom_intrinsics/serial/impl.cpp index 9b7c6b51..aba55cd2 100644 --- a/tests/functional/data/misc/custom_intrinsics/openmp/impl.h +++ b/tests/functional/data/misc/custom_intrinsics/serial/impl.cpp @@ -1,5 +1,3 @@ -#pragma once - //INFO: needed for std::isnan #include diff --git a/tests/functional/data/transpiler/backends/cuda/user_intrinsics/user_intrinsics_ref.cpp b/tests/functional/data/transpiler/backends/cuda/user_intrinsics/user_intrinsics_ref.cpp index 4e658c1a..e12a2832 100644 --- a/tests/functional/data/transpiler/backends/cuda/user_intrinsics/user_intrinsics_ref.cpp +++ b/tests/functional/data/transpiler/backends/cuda/user_intrinsics/user_intrinsics_ref.cpp @@ -1,6 +1,7 @@ #include // INFO: transpiling to cuda backend already include necessary header + __device__ bool okl_is_nan(float value) { return isnan(value) != 0; } extern "C" __global__ __launch_bounds__(32) void _occa_zero_nans_0(float *vec) { diff --git a/tests/functional/data/transpiler/backends/dpcpp/user_intrinsics/user_intrinsics_ref.cpp b/tests/functional/data/transpiler/backends/dpcpp/user_intrinsics/user_intrinsics_ref.cpp index 218cb5a9..4d6a1cf7 100644 --- a/tests/functional/data/transpiler/backends/dpcpp/user_intrinsics/user_intrinsics_ref.cpp +++ b/tests/functional/data/transpiler/backends/dpcpp/user_intrinsics/user_intrinsics_ref.cpp @@ -1,5 +1,4 @@ #include -using namespace sycl; // INFO: from documentation // isNaN @@ -10,6 +9,8 @@ SYCL_EXTERNAL bool okl_is_nan(float value) { return __binary32_isNaN(value) == 1; } +using namespace sycl; + extern "C" [[sycl::reqd_work_group_size(1, 1, 32)]] void _occa_zero_nans_0(sycl::queue *queue_, sycl::nd_range<3> *range_, float *vec) { queue_->submit([&](sycl::handler &handler_) { diff --git a/tests/functional/data/transpiler/backends/hip/user_intrinsics/user_intrinsics_ref.cpp b/tests/functional/data/transpiler/backends/hip/user_intrinsics/user_intrinsics_ref.cpp index 1c59dcfb..fb0989ac 100644 --- a/tests/functional/data/transpiler/backends/hip/user_intrinsics/user_intrinsics_ref.cpp +++ b/tests/functional/data/transpiler/backends/hip/user_intrinsics/user_intrinsics_ref.cpp @@ -1,7 +1,8 @@ #include // INFO: relate to the documentation is must be supported natively -// https://rocm.docs.amd.com/projects/HIP/en/latest/reference/kernel_language.html#single-precision-mathematical-functions +// https://rocm.docs.amd.com/projects/HIP/en/latest/reference/kernel_language.html#single-precision-mathematical-functions + __device__ bool okl_is_nan(float value) { return isnan(value); } extern "C" __global__ __launch_bounds__(32) void _occa_zero_nans_0(float *vec) { diff --git a/tests/functional/data/transpiler/backends/openmp/user_intrinsics/user_intrinsics_ref.cpp b/tests/functional/data/transpiler/backends/openmp/user_intrinsics/user_intrinsics_ref.cpp index 21a4bc43..6922b06a 100644 --- a/tests/functional/data/transpiler/backends/openmp/user_intrinsics/user_intrinsics_ref.cpp +++ b/tests/functional/data/transpiler/backends/openmp/user_intrinsics/user_intrinsics_ref.cpp @@ -1,7 +1,9 @@ +// INFO: needed for std::isnan #include -// INFO: needed for std::isnan -bool okl_is_nan(float value) { return std::isnan(value); } +bool okl_is_nan(float value) { + return std::isnan(value); +} extern "C" void zero_nans(float *vec) { #pragma omp parallel for diff --git a/tests/functional/data/transpiler/backends/serial/user_intrinsics/user_intrinsics_ref.cpp b/tests/functional/data/transpiler/backends/serial/user_intrinsics/user_intrinsics_ref.cpp index 234cac41..2c10dc0e 100644 --- a/tests/functional/data/transpiler/backends/serial/user_intrinsics/user_intrinsics_ref.cpp +++ b/tests/functional/data/transpiler/backends/serial/user_intrinsics/user_intrinsics_ref.cpp @@ -1,7 +1,9 @@ +//INFO: needed for std::isnan #include -// INFO: needed for std::isnan -bool okl_is_nan(float value) { return std::isnan(value); } +bool okl_is_nan(float value) { + return std::isnan(value); +} extern "C" void zero_nans(float *vec) { for (int i = 0; i < 32; ++i) {