forked from LuisaGroup/LuisaCompute
-
Notifications
You must be signed in to change notification settings - Fork 0
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
Merge remote-tracking branch 'origin/dev-fallback-backend' into dev-xir
- Loading branch information
Showing
112 changed files
with
17,576 additions
and
11,706 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -1,54 +1,44 @@ | ||
| #pragma once | ||
|
|
||
| #if defined(__x86_64__) || defined(_M_X64) | ||
| #define LUISA_ARCH_X86_64 | ||
| #elif defined(__aarch64__) | ||
| #define LUISA_ARCH_ARM64 | ||
| #else | ||
| #error Unsupported architecture | ||
| #endif | ||
|
|
||
| #if defined(LUISA_ARCH_X86_64) | ||
|
|
||
| #include <immintrin.h> | ||
| #include <cstdint> | ||
| #include <cassert> | ||
|
|
||
| #define LUISA_INTRIN_PAUSE() _mm_pause() | ||
|
|
||
| namespace luisa { | ||
| using float16_t = int16_t; | ||
| using float32x4_t = __m128; | ||
| }// namespace luisa | ||
|
|
||
| #elif defined(LUISA_ARCH_ARM64) | ||
|
|
||
| #include <arm_neon.h> | ||
|
|
||
| namespace luisa { | ||
| using float16_t = ::float16_t; | ||
| using float32x4_t = ::float32x4_t; | ||
| }// namespace luisa | ||
|
|
||
| #define LUISA_INTRIN_PAUSE() asm volatile("isb") | ||
|
|
||
| #else | ||
|
|
||
| #include <thread> | ||
| #define LUISA_INTRIN_PAUSE() std::this_thread::yield() | ||
|
|
||
| #endif | ||
|
|
||
| ////////////// assume | ||
| #ifdef NDEBUG // assume only enabled in non-debug mode. | ||
| #if defined(__clang__)// Clang | ||
| #define LUISA_ASSUME(x) (__builtin_assume(x)) | ||
| #elif defined(_MSC_VER)// MSVC | ||
| #define LUISA_ASSUME(x) (__assume(x)) | ||
| #else// GCC | ||
| #define LUISA_ASSUME(x) \ | ||
| if (!(x)) __builtin_unreachable() | ||
| #endif | ||
| #else | ||
| #define LUISA_ASSUME(expression) assert(expression) | ||
| #endif | ||
| #pragma once | ||
|
|
||
| #if defined(__x86_64__) || defined(_M_X64) | ||
| #define LUISA_ARCH_X86_64 | ||
| #elif defined(__aarch64__) | ||
| #define LUISA_ARCH_ARM64 | ||
| #else | ||
| #error Unsupported architecture | ||
| #endif | ||
|
|
||
| #if defined(LUISA_ARCH_X86_64) | ||
|
|
||
| #include <immintrin.h> | ||
| #include <cstdint> | ||
| #include <cassert> | ||
|
|
||
| #define LUISA_INTRIN_PAUSE() _mm_pause() | ||
|
|
||
| #elif defined(LUISA_ARCH_ARM64) | ||
|
|
||
| #include <arm_neon.h> | ||
|
|
||
| #define LUISA_INTRIN_PAUSE() asm volatile("isb") | ||
|
|
||
| #else | ||
|
|
||
| #include <thread> | ||
| #define LUISA_INTRIN_PAUSE() std::this_thread::yield() | ||
|
|
||
| #endif | ||
|
|
||
| ////////////// assume | ||
| #ifdef NDEBUG // assume only enabled in non-debug mode. | ||
| #if defined(__clang__)// Clang | ||
| #define LUISA_ASSUME(x) (__builtin_assume(x)) | ||
| #elif defined(_MSC_VER)// MSVC | ||
| #define LUISA_ASSUME(x) (__assume(x)) | ||
| #else// GCC | ||
| #define LUISA_ASSUME(x) \ | ||
| if (!(x)) __builtin_unreachable() | ||
| #endif | ||
| #else | ||
| #define LUISA_ASSUME(expression) assert(expression) | ||
| #endif |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -1,67 +1,69 @@ | ||
| #pragma once | ||
|
|
||
| #include <luisa/core/concepts.h> | ||
| #include <luisa/core/stl/optional.h> | ||
| #include <luisa/core/stl/functional.h> | ||
| #include <luisa/runtime/rhi/command.h> | ||
|
|
||
| #ifdef LUISA_ENABLE_API | ||
| #include <luisa/api/common.h> | ||
| #endif | ||
| namespace lc::validation { | ||
| class Device; | ||
| }// namespace lc::validation | ||
| namespace luisa::compute { | ||
|
|
||
| class LC_RUNTIME_API CommandList : concepts::Noncopyable { | ||
| friend class lc::validation::Device; | ||
|
|
||
| public: | ||
| class Commit; | ||
| using CommandContainer = luisa::vector<luisa::unique_ptr<Command>>; | ||
| using CallbackContainer = luisa::vector<luisa::move_only_function<void()>>; | ||
|
|
||
| private: | ||
| CommandContainer _commands; | ||
| CallbackContainer _callbacks; | ||
| bool _committed{false}; | ||
|
|
||
| public: | ||
| CommandList() noexcept = default; | ||
| ~CommandList() noexcept; | ||
| CommandList(CommandList &&another) noexcept; | ||
| CommandList &operator=(CommandList &&rhs) noexcept = delete; | ||
| [[nodiscard]] static CommandList create(size_t reserved_command_size = 0u, | ||
| size_t reserved_callback_size = 0u) noexcept; | ||
|
|
||
| void reserve(size_t command_size, size_t callback_size) noexcept; | ||
| CommandList &operator<<(luisa::unique_ptr<Command> &&cmd) noexcept; | ||
| CommandList &append(luisa::unique_ptr<Command> &&cmd) noexcept; | ||
| CommandList &add_callback(luisa::move_only_function<void()> &&callback) noexcept; | ||
| void clear() noexcept; | ||
| [[nodiscard]] auto commands() const noexcept { return luisa::span{_commands}; } | ||
| [[nodiscard]] auto callbacks() const noexcept { return luisa::span{_callbacks}; } | ||
| [[nodiscard]] CommandContainer steal_commands() noexcept; | ||
| [[nodiscard]] CallbackContainer steal_callbacks() noexcept; | ||
| [[nodiscard]] auto empty() const noexcept { return _commands.empty() && _callbacks.empty(); } | ||
| [[nodiscard]] Commit commit() noexcept; | ||
| }; | ||
|
|
||
| class CommandList::Commit { | ||
|
|
||
| private: | ||
| CommandList _list; | ||
|
|
||
| private: | ||
| friend class CommandList; | ||
| explicit Commit(CommandList &&list) noexcept | ||
| : _list{std::move(list)} {} | ||
| Commit(Commit &&) noexcept = default; | ||
|
|
||
| public: | ||
| Commit &operator=(Commit &&) noexcept = delete; | ||
| Commit &operator=(const Commit &) noexcept = delete; | ||
| [[nodiscard]] auto command_list() && noexcept { return std::move(_list); } | ||
| }; | ||
|
|
||
| }// namespace luisa::compute | ||
| #pragma once | ||
|
|
||
| #include <luisa/core/concepts.h> | ||
| #include <luisa/core/stl/optional.h> | ||
| #include <luisa/core/stl/functional.h> | ||
| #include <luisa/runtime/rhi/command.h> | ||
|
|
||
| #ifdef LUISA_ENABLE_API | ||
| #include <luisa/api/common.h> | ||
| #endif | ||
|
|
||
| namespace lc::validation { | ||
| class Device; | ||
| }// namespace lc::validation | ||
|
|
||
| namespace luisa::compute { | ||
|
|
||
| class LC_RUNTIME_API CommandList : concepts::Noncopyable { | ||
| friend class lc::validation::Device; | ||
|
|
||
| public: | ||
| class Commit; | ||
| using CommandContainer = luisa::vector<luisa::unique_ptr<Command>>; | ||
| using CallbackContainer = luisa::vector<luisa::move_only_function<void()>>; | ||
|
|
||
| private: | ||
| CommandContainer _commands; | ||
| CallbackContainer _callbacks; | ||
| bool _committed{false}; | ||
|
|
||
| public: | ||
| CommandList() noexcept = default; | ||
| ~CommandList() noexcept; | ||
| CommandList(CommandList &&another) noexcept; | ||
| CommandList &operator=(CommandList &&rhs) noexcept = delete; | ||
| [[nodiscard]] static CommandList create(size_t reserved_command_size = 0u, | ||
| size_t reserved_callback_size = 0u) noexcept; | ||
|
|
||
| void reserve(size_t command_size, size_t callback_size) noexcept; | ||
| CommandList &operator<<(luisa::unique_ptr<Command> &&cmd) noexcept; | ||
| CommandList &append(luisa::unique_ptr<Command> &&cmd) noexcept; | ||
| CommandList &add_callback(luisa::move_only_function<void()> &&callback) noexcept; | ||
| void clear() noexcept; | ||
| [[nodiscard]] auto commands() const noexcept { return luisa::span{_commands}; } | ||
| [[nodiscard]] auto callbacks() const noexcept { return luisa::span{_callbacks}; } | ||
| [[nodiscard]] CommandContainer steal_commands() noexcept; | ||
| [[nodiscard]] CallbackContainer steal_callbacks() noexcept; | ||
| [[nodiscard]] auto empty() const noexcept { return _commands.empty() && _callbacks.empty(); } | ||
| [[nodiscard]] Commit commit() noexcept; | ||
| }; | ||
|
|
||
| class CommandList::Commit { | ||
|
|
||
| private: | ||
| CommandList _list; | ||
|
|
||
| private: | ||
| friend class CommandList; | ||
| explicit Commit(CommandList &&list) noexcept | ||
| : _list{std::move(list)} {} | ||
| Commit(Commit &&) noexcept = default; | ||
|
|
||
| public: | ||
| Commit &operator=(Commit &&) noexcept = delete; | ||
| Commit &operator=(const Commit &) noexcept = delete; | ||
| [[nodiscard]] auto command_list() && noexcept { return std::move(_list); } | ||
| }; | ||
|
|
||
| }// namespace luisa::compute |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,19 @@ | ||
| #pragma once | ||
|
|
||
| #include <luisa/core/stl/unordered_map.h> | ||
| #include <luisa/xir/module.h> | ||
| #include <luisa/xir/instructions/outline.h> | ||
|
|
||
| namespace luisa::compute::xir { | ||
|
|
||
| // This pass will outline all outline instructions in the module. | ||
| // Information about the outlined functions will be returned. | ||
|
|
||
| struct OutlineInfo { | ||
| luisa::unordered_map<OutlineInst *, Function *> outlines; | ||
| }; | ||
|
|
||
| LC_XIR_API OutlineInfo outline_pass_run_on_function(Module *module, Function *function) noexcept; | ||
| LC_XIR_API OutlineInfo outline_pass_run_on_module(Module *module) noexcept; | ||
|
|
||
| }// namespace luisa::compute::xir |
Oops, something went wrong.