diff --git a/README.md b/README.md index a78ee0e9..0ed9d602 100644 --- a/README.md +++ b/README.md @@ -56,11 +56,11 @@ FPS comparison of rendering [Franka](https://github.com/Genesis-Embodied-AI/Gene Resolution: 128x128

- FPS of gs-madrona rasterizer vs raytracer + FPS of gs-madrona rasterizer vs raytracer

- FPS of gs-madrona rasterizer vs raytracer + FPS of gs-madrona rasterizer vs raytracer

## Install (Linux Only) @@ -104,22 +104,6 @@ renderer = gs.options.renderers.BatchRenderer( ) ``` -### Performance Benchmark -For comprehensive performance benchmarking across multiple renderers (Madrona, Omniverse, PyRender, ManiSkill), please refer to the detailed documentation in `scripts/perf_benchmark/README.md`. - -The benchmark suite includes: -- Multi-renderer performance testing -- Batch size and resolution scaling tests -- Rasterizer vs raytracer comparisons -- Automated report generation -- Asset preprocessing utilities - -Quick start: -```bash -cd scripts/perf_benchmark -python batch_benchmark.py -f benchmark_config_smoke_test.yml -``` - ## Acknowledgments The development of gs-madrona is actively supported by [Genesis AI](https://genesis-ai.company/). diff --git a/docs/perf_franka_128x128_plot.png b/docs/perf_franka_128x128_plot.png new file mode 100644 index 00000000..9d3e9e77 Binary files /dev/null and b/docs/perf_franka_128x128_plot.png differ diff --git a/docs/perf_franka_128x128_table.png b/docs/perf_franka_128x128_table.png new file mode 100644 index 00000000..eb00e7b8 Binary files /dev/null and b/docs/perf_franka_128x128_table.png differ diff --git a/include/madrona/context.inl b/include/madrona/context.inl index 153db67e..c62658da 100644 --- a/include/madrona/context.inl +++ b/include/madrona/context.inl @@ -112,12 +112,6 @@ inline void Context::iterateQuery(const Query &query, Fn &&fn) std::forward(fn)); } -#ifdef MADRONA_USE_JOB_SYSTEM -JobID Context::currentJobID() const -{ - return cur_job_id_; -} -#endif #ifdef MADRONA_MW_MODE WorldID Context::worldID() const @@ -126,267 +120,5 @@ WorldID Context::worldID() const } #endif -#if 0 - -class Context { - AllocContext mem; - - template - inline ArchetypeRef archetype(); - - template - inline void clearArchetype(); - - template - inline void clearTemporaries(); - - template - inline Query query(); - - template - inline void forEach(const Query &query, Fn &&fn); - - template - inline uint32_t numMatches(const Query &query); - - // Jobs - template - inline JobID submit(Fn &&fn, bool is_child = true, - DepTs && ... dependencies); - - template - inline JobID submitN(Fn &&fn, uint32_t num_invocations, - bool is_child = true, - DepTs && ... dependencies); - - // FIXME: currently this function requires that the query reference - // is valid at least until the returned job is completed. - template - inline JobID parallelFor(const Query &query, Fn &&fn, - bool is_child = true, - DepTs && ... dependencies); - - template - inline JobID ioRead(const char *path, Fn &&fn, bool is_child = true, - DepTs && ... dependencies); - -#ifdef MADRONA_USE_JOB_SYSTEM - inline JobID currentJobID() const; -#endif - - template - inline JobID submitImpl(Fn &&fn, bool is_child, DepTs && ... dependencies); - - template - inline JobID submitNImpl(Fn &&fn, uint32_t num_invocations, bool is_child, - DepTs && ... dependencies); - - template - inline JobID parallelForImpl(const Query &query, Fn &&fn, - bool is_child, DepTs && ... dependencies); - -private: - template - inline JobID submitNImpl(Fn &&fn, uint32_t num_invocations, JobID parent_id, - DepTs && ... dependencies); - -#ifdef MADRONA_USE_JOB_SYSTEM - JobManager * const job_mgr_; - StateManager * const state_mgr_; - StateCache * const state_cache_; - IOManager * const io_mgr_; - const int worker_idx_; - JobID cur_job_id_; -#endif - -friend class JobManager; -}; - -template -ArchetypeRef Context::archetype() -{ - return state_mgr_->archetype( - MADRONA_MW_COND(cur_world_id_)); -} - -template -void Context::clearArchetype() -{ - state_mgr_->clear(MADRONA_MW_COND(cur_world_id_,) - *state_cache_, false); -} - -template -void Context::clearTemporaries() -{ -} - -template -Query Context::query() -{ - return state_mgr_->query(); -} - -template -void Context::forEach(const Query &query, Fn &&fn) -{ - state_mgr_->iterateQuery(MADRONA_MW_COND(cur_world_id_,) query, - std::forward(fn)); -} - -template -uint32_t Context::numMatches(const Query &query) -{ - uint32_t num_entities = 0; - state_mgr_->iterateArchetypes(MADRONA_MW_COND(cur_world_id_,) query, - [&](int num_rows, auto ...) { - num_entities += num_rows; - }); - - return num_entities; -} - -template -JobID Context::submit(Fn &&fn, bool is_child, Deps && ... dependencies) -{ - return submitImpl(std::forward(fn), is_child, - std::forward(dependencies)...); -} - -template -JobID Context::submitN(Fn &&fn, uint32_t num_invocations, - bool is_child, Deps && ... dependencies) -{ - return submitNImpl(std::forward(fn), num_invocations, - is_child, std::forward(dependencies)...); -} - -template -JobID Context::parallelFor(const Query &query, Fn &&fn, - bool is_child, Deps && ... dependencies) -{ - return parallelForImpl(query, std::forward(fn), is_child, - std::forward(dependencies)...); -} - - -template -inline JobID Context::ioRead(const char *path, Fn &&fn, - bool is_child, Deps && ... dependencies) -{ - IOPromise promise = io_mgr_->makePromise(); - Job job = makeJob([promise, fn=std::move(fn), io_mgr=io_mgr_]( - Context &ctx) { - fn(ctx, io_mgr->getBuffer(promise)); - }); - - io_mgr_->load(promise, path, job); - - (void)is_child; - ( (void)dependencies, ... ); - - return JobID::none(); -} - -// FIXME: implement is_child, dependencies, num_invocations -template -JobID Context::submitImpl(Fn &&fn, bool is_child, - Deps &&... dependencies) -{ - JobID parent_id = is_child ? cur_job_id_ : JobID::none(); - - return job_mgr_->queueJob(worker_idx_, - std::forward(fn), 0, parent_id, - MADRONA_MW_COND(cur_world_id_, ) JobPriority::Normal, - std::forward(dependencies)...); -} - -template -JobID Context::submitNImpl(Fn &&fn, uint32_t num_invocations, bool is_child, - Deps && ...dependencies) -{ - assert(num_invocations > 0); - JobID parent_id = is_child ? cur_job_id_ : JobID::none(); - - return submitNImpl(std::forward(fn), num_invocations, - parent_id, std::forward(dependencies)...); -} - -template -JobID Context::parallelForImpl(const Query &query, Fn &&fn, - bool is_child, Deps && ... dependencies) -{ - if (query.numMatchingArchetypes() == 0) { - return JobID::none(); - } - - // FIXME: add isRunnable check in addition to no dependencies - if constexpr (sizeof...(dependencies) == 0) { - JobID parent_id = is_child ? cur_job_id_ : JobID::none(); - - // Additional optimization: skip this proxy ID when only 1 archetype - // is present (in fact for > 1 archetype might make sense to just - // use the else codepath). - JobID proxy_id = job_mgr_->reserveProxyJobID(worker_idx_, parent_id); - - state_mgr_->iterateArchetypes(MADRONA_MW_COND(cur_world_id_,) - query, [this, &fn, proxy_id](int num_rows, auto ...ptrs) { - if (num_rows == 0) { - return; - } - - // Clang complains this is unused without this-> - this->submitNImpl( - [fn = Fn(fn), ptrs...](ContextT &ctx, uint32_t idx) { - fn(ctx, ptrs[idx]...); - }, num_rows, proxy_id); - }); - - // Note that even though we "relinquish" the id here, it is still safe - // to return the ID, since the generation stored in the ID will simply - // be invalid if the entire parallelFor job finishes, just like a normal - // job id. - job_mgr_->relinquishProxyJobID(worker_idx_, proxy_id); - - return proxy_id; - } else { - return submitImpl([fn = std::forward(fn), &query] ( - ContextT &ctx) { - ctx.state_mgr_->iterateArchetypes( - MADRONA_MW_COND(ctx.cur_world_id_,) query, - [&ctx, &fn](int num_rows, auto ...ptrs) { - if (num_rows == 0) { - return; - } - - // FIXME reconsider copying ptrs into the closure here - // FIXME currently copies the user function's closure - // Could allow making a fake jobs with data but not a function - // by extending reserveProxyJobID - that job could be dependent - // on the parallel for job and hold the user function closure. - // If we allowed runtime determined # of dependencies, the - // fast path (no dependencies above) could return the dependent - // data-only job rather than using the fake ID as a parent - ctx.template submitNImpl( - [fn = Fn(fn), ptrs...](ContextT &ctx, uint32_t idx) { - fn(ctx, ptrs[idx]...); - }, num_rows, true); - }); - }, is_child, dependencies...); - } -} - -template -JobID Context::submitNImpl(Fn &&fn, uint32_t num_invocations, JobID parent_id, - Deps && ...dependencies) -{ - return job_mgr_->queueJob(worker_idx_, - std::forward(fn), num_invocations, parent_id, - MADRONA_MW_COND(cur_world_id_, ) JobPriority::Normal, - std::forward(dependencies)...); -} -#endif } diff --git a/include/madrona/custom_context.inl b/include/madrona/custom_context.inl index 09ad1285..10f35d30 100644 --- a/include/madrona/custom_context.inl +++ b/include/madrona/custom_context.inl @@ -15,55 +15,5 @@ CustomContext::CustomContext(DataT *world_data, : Context(world_data, worker_init) {} -#if 0 - -template -class CustomContext : public Context { -public: - template - inline JobID submit(Fn &&fn, bool is_child = true, - Deps && ... dependencies); - - template - inline JobID submitN(Fn &&fn, uint32_t num_invocations, - bool is_child = true, - Deps && ... dependencies); - - template - inline JobID parallelFor(const Query &query, Fn &&fn, - bool is_child = true, - Deps && ... dependencies); -} - -template -template -JobID CustomContext::submit(Fn &&fn, bool is_child, - Deps && ... dependencies) -{ - return submitImpl(std::forward(fn), is_child, - std::forward(dependencies)...); -} - -template -template -JobID CustomContext::submitN(Fn &&fn, - uint32_t num_invocations, bool is_child, Deps && ... dependencies) -{ - return submitNImpl( - std::forward(fn), num_invocations, is_child, - std::forward(dependencies)...); -} - -template -template -JobID CustomContext::parallelFor( - const Query &query, Fn &&fn, bool is_child, - Deps && ... dependencies) -{ - return parallelForImpl(query, std::forward(fn), is_child, - std::forward(dependencies)...); -} - -#endif } diff --git a/src/core/context.cpp b/src/core/context.cpp index 1cd4190d..b005b7a8 100644 --- a/src/core/context.cpp +++ b/src/core/context.cpp @@ -13,14 +13,6 @@ namespace madrona { Context::Context(WorldBase *world_data, const WorkerInit &init) : data_(world_data) -#ifdef MADRONA_USE_JOB_SYSTEM - , job_mgr_(init.jobMgr), - state_mgr_(init.stateMgr), - state_cache_(init.stateCache), - io_mgr_(nullptr), - worker_idx_(init.workerIdx), - cur_job_id_(JobID::none()) -#endif #ifdef MADRONA_USE_TASK_GRAPH , state_mgr_(init.stateMgr) , state_cache_(init.stateCache) diff --git a/src/core/worker_init.hpp b/src/core/worker_init.hpp index 1b54ade5..c25a1b91 100644 --- a/src/core/worker_init.hpp +++ b/src/core/worker_init.hpp @@ -14,12 +14,6 @@ namespace madrona { struct WorkerInit { -#ifdef MADRONA_USE_JOB_SYSTEM - JobManager *jobMgr; - StateManager *stateMgr; - StateCache *stateCache; - int workerIdx; -#endif #ifdef MADRONA_USE_TASK_GRAPH StateManager *stateMgr; StateCache *stateCache; diff --git a/src/mw/device/include/madrona/context.inl b/src/mw/device/include/madrona/context.inl index 6618bcb8..0de45efa 100644 --- a/src/mw/device/include/madrona/context.inl +++ b/src/mw/device/include/madrona/context.inl @@ -126,205 +126,5 @@ inline void Context::iterateQuery(Query &query, Fn&& fn) }); } -#if 0 - -class Context { - template - inline JobID submit(Fn &&fn, bool is_child = true, - DepTs && ...dependencies); - - template - inline JobID submitN(Fn &&fn, uint32_t num_invocations, - bool is_child = true, DepTs && ...dependencies); - - template - inline JobID parallelFor(const Query &query, Fn &&fn, - bool is_child = true, DepTs && ... dependencies); - -#if 0 - template - inline JobID submitImpl(Fn &&fn, bool is_child, - DepTs && ... dependencies); - - template - inline JobID submitNImpl(Fn &&fn, uint32_t num_invocations, bool is_child, - DepTs && ... dependencies); - - template - inline JobID parallelForImpl(const Query &query, Fn &&fn, - bool is_child, DepTs && ... dependencies); -#endif - - void markJobFinished(); - - inline JobID currentJobID() const { return job_id_; } - - inline StateManager & state(); - - WaveInfo computeWaveInfo(); - - JobID waveSetupNewJob(uint32_t func_id, bool link_parent, - uint32_t num_invocations, uint32_t bytes_per_job, - void **thread_data_store); - - JobContainerBase * allocJob(uint32_t bytes_per_job, WaveInfo wave_info); - - inline void stageChildJob(uint32_t func_id, uint32_t num_combined_jobs, - uint32_t bytes_per_job, void *containers); - - JobID job_id_; - - uint32_t lane_id_; -} - -namespace mwGPU { - -// This function is executed at the thread-block granularity. -// num_invocations <= consts::numMegakernelThreads -template -__attribute__((used, always_inline)) -inline void jobEntry(JobContainerBase *job_data, - uint32_t *data_indices, - uint32_t *invocation_offsets, - uint32_t num_invocations, - uint32_t grid_id) -{ - uint32_t lane_id = threadIdx.x % consts::numWarpThreads; - - uint32_t invocation_idx = - threadIdx.x + blockIdx.x * consts::numJobLaunchKernelThreads; - - if (invocation_idx >= num_invocations) { - return; - } - - uint32_t data_idx = data_indices[invocation_idx]; - uint32_t local_offset = invocation_offsets[invocation_idx]; - - static_assert(std::is_trivially_destructible_v); - - ContainerT &job_container = - static_cast(job_data)[data_idx]; - - ContextT ctx = JobManager::makeContext(job_container.jobID, - grid_id, job_container.worldID, lane_id); - - (job_container.fn)(ctx, local_offset); - - ctx.markJobFinished(num_invocations); -} - -template -struct JobFuncIDBase { - static uint32_t id; -}; - -template ) = - jobEntry> -struct JobFuncID : JobFuncIDBase {}; - -} - -Context::Context(WorldBase *world_data, WorkerInit &&init) - : data_(world_data), - job_id_(init.jobID), - grid_id_(init.gridID), - world_id_(init.worldID), - lane_id_(init.laneID) -{} - -StateManager & Context::state() -{ - return *(StateManager *)mwGPU::GPUImplConsts::get().stateManagerAddr; -} - -template -JobID Context::submit(Fn &&fn, bool is_child, DepTs && ... dependencies) -{ - return submitImpl(std::forward(fn), 1, is_child, - std::forward(dependencies)...); -} - -template -JobID Context::submitN(Fn &&fn, uint32_t num_invocations, - bool is_child, DepTs && ... dependencies) -{ - return submitImpl(std::forward(fn), num_invocations, is_child, - std::forward(dependencies)...); -} - -template -JobID Context::parallelFor(const Query &query, Fn &&fn, - bool is_child, DepTs && ... dependencies) -{ - return parallelForImpl(query, std::forward(fn), is_child, - std::forward(dependencies)...); -} - - -template -JobID Context::submitImpl(Fn &&fn, bool is_child, DepTs && ... dependencies) -{ - auto wrapper = [fn = std::forward(fn)](ContextT &ctx, uint32_t) { - fn(ctx); - }; - - return submitNImpl(std::move(wrapper), 1, is_child, - std::forward(dependencies)...); -} - -template -JobID Context::submitNImpl(Fn &&fn, uint32_t num_invocations, bool is_child, - DepTs &&... dependencies) -{ - using namespace mwGPU; - - constexpr std::size_t num_deps = sizeof...(DepTs); - - using ContainerT = JobContainer; - - UserJobTracker *user_trackers = - JobManager::get()->getUserJobTrackers(); - - auto wave_info = computeWaveInfo(); - - // Consolidate dependencies across warp - // num_deps is guaranteed to be the same across activemask here -#if 0 - uint32_t num_unique_dependencies = 0; - auto countUniqueDeps = [&](JobID dep) { - UserJobTracker &user_tracker = user_trackers[dep.id]; - - uint32_t shared_id = dep.gen == user_tracker.gen ? - user_tracker.sharedID : ~0_u32; - - uint32_t match_mask = __match_any_sync(wave_info.activeMask, - shared_id); - - uint32_t top_match = getHighestSetBit(match_mask); - - uint32_t num_add = (top_match == lane_id_ && shared_id != ~0_u32) ? - 1 : 0; - num_unique_dependencies += - __reduce_add_sync(wave_info.activeMask, num_add); - }; -#endif - - uint32_t func_id = mwGPU::JobFuncID::id; - - void *thread_data_store; - JobID job_id = waveSetupNewJob(func_id, is_child, num_invocations, - sizeof(ContainerT), &thread_data_store); - - new (thread_data_store) ContainerT(job_id, world_id_, - num_invocations, - std::forward(fn), - std::forward(dependencies)...); - - return queue_job_id; -} -#endif } diff --git a/src/render/CMakeLists.txt b/src/render/CMakeLists.txt index 73552b37..0c4016b4 100644 --- a/src/render/CMakeLists.txt +++ b/src/render/CMakeLists.txt @@ -22,7 +22,6 @@ add_library(madrona_render_core STATIC render_ctx.hpp render_ctx.cpp batch_renderer.hpp batch_renderer.cpp render_common.hpp - image_util.cpp ) target_compile_definitions(madrona_render_core PUBLIC diff --git a/src/render/image_util.cpp b/src/render/image_util.cpp deleted file mode 100644 index f25fa1b3..00000000 --- a/src/render/image_util.cpp +++ /dev/null @@ -1,2 +0,0 @@ -#define STB_IMAGE_WRITE_IMPLEMENTATION -#include