Pad element count to next power of two for bitonic correctness. Out-of-range threads early-exit in the shader via element_count.
Total bitonic passes: sum of (stage+1) for stage in [0, stage_count).
Recover (stage, step) from the flat pass index. Passes enumerate: stage=0 step=0, stage=1 step=0,1, stage=2 step=0,1,2, ...
619{
620 const auto feature_key = this->template get_parameter_or_default<std::string>("feature_key", std::string("feature"));
621 const bool ascending = this->template get_parameter_or_default<bool>("ascending", true);
622 const auto gpu_thresh = this->template get_parameter_or_default<uint32_t>("gpu_sort_threshold", 0U);
623 const auto n =
static_cast<uint32_t
>(
input.data.regions.size());
624
625 Datum<Kakshya::RegionGroup> out =
input;
626
627 if (n == 0)
628 return out;
629
630 if (gpu_thresh > 0 && n >= gpu_thresh) {
631 struct SortPC {
632 uint32_t element_count;
633 uint32_t step;
634 uint32_t stage;
635 uint32_t ascending;
636 };
637
638 std::vector<float> values(n);
639 for (uint32_t i = 0; i < n; ++i) {
640 auto attr = out.data.regions[i].get_attribute<double>(feature_key);
641 values[i] = attr ? static_cast<float>(*attr) : 0.0F;
642 }
643
644 std::vector<uint32_t> indices(n);
645 std::iota(indices.begin(), indices.end(), 0U);
646
647
648
649 uint32_t n_padded = 1U;
650 while (n_padded < n)
651 n_padded <<= 1U;
652
653
654 uint32_t stage_count = 0
U;
655 {
656 uint32_t tmp = n_padded;
657 while (tmp > 1U) {
658 ++stage_count;
659 tmp >>= 1U;
660 }
661 }
662 const uint32_t n_passes = stage_count * (stage_count + 1U) / 2U;
663
664 auto executor = std::make_shared<ShaderExecutionContext<>>(
667 .workgroup_size = { 256, 1, 1 },
668 .push_constant_size = sizeof(SortPC) });
669
672 .set_output_size(1, n * sizeof(uint32_t));
673
674 executor->set_multipass(n_passes,
675 [n, ascending](uint32_t
pass,
void* pc_data) {
676
677
678
679
681 uint32_t remaining =
pass;
682 while (remaining >= stage + 1U) {
683 remaining -= stage + 1U;
684 ++stage;
685 }
686 const uint32_t step = stage - remaining;
687
688 const SortPC pc {
689 .element_count = n,
690 .step = step,
691 .stage = stage,
692 .ascending = ascending ? 1U : 0
U,
693 };
694 std::memcpy(pc_data, &pc, sizeof(SortPC));
695 });
696
697 auto gpu_sorter = std::make_shared<GpuSorter<>>(executor);
698
699 Datum<std::vector<Kakshya::DataVariant>> sort_input {
701 };
702
703 auto result = gpu_sorter->apply_operation(sort_input);
704 const auto sorted_indices = ShaderExecutionContext<>::read_output<uint32_t>(result, 1);
705
706 std::vector<Kakshya::Region> reordered(n);
707 for (uint32_t i = 0; i < n; ++i)
708 reordered[i] = out.data.regions[sorted_indices[i]];
709
710 out.data.regions = std::move(reordered);
711 out.data.current_region_index = 0;
712 out.data.active_indices.clear();
713 return out;
714 }
715
716 out.data.sort_by_attribute(feature_key);
717
718 if (!ascending)
719 std::ranges::reverse(out.data.regions);
720
721 out.data.current_region_index = 0;
722 out.data.active_indices.clear();
723
724 return out;
725}
Core::GlobalInputConfig input
std::variant< std::vector< double >, std::vector< float >, std::vector< uint8_t >, std::vector< uint16_t >, std::vector< uint32_t >, std::vector< std::complex< float > >, std::vector< std::complex< double > >, std::vector< glm::vec2 >, std::vector< glm::vec3 >, std::vector< glm::vec4 >, std::vector< glm::mat4 > > DataVariant
Multi-type data storage for different precision needs.
Portal::Graphics::GpuComputeConfig GpuComputeConfig
enum MayaFlux::Portal::Graphics::GpuBufferBinding::ElementType FLOAT32