Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Reduce the number of kernels to compile in merge-sort #1872

Open
wants to merge 4 commits into
base: main
Choose a base branch
from

Conversation

dmitriy-sobolev
Copy link
Contributor

@dmitriy-sobolev dmitriy-sobolev commented Sep 23, 2024

This PR implements a TODO:

    // TODO: split the submitter into multiple ones to avoid extra compilation of kernels:
    // - _LeafSortKernel and _CopyBackKernel do not need _IndexT

It results in reduced kernel compilation times. The table below shows a JIT compilation speed up for different devices:

  Named Kernels: JIT speedup, times Unnamed Kernels: JIT speedup, times
Intel® Data Center GPU Max Series (zeModuleCreate calls) 1.18 1.61
Intel(R) Xeon(R) Platinum 8480+ (clBuildProgram calls) 1.26 1.85

There is no observable impact on the execution time besides kernel compilation.

I used the following program as a benchmark:

#include <sycl/sycl.hpp>

#include <oneapi/dpl/algorithm>
#include <oneapi/dpl/execution>

#include <iostream>
#include <algorithm>

int main()
{
    using T = int;
    const int n = 50'000;
    sycl::queue q;
    T* data = sycl::malloc_shared<T>(n, q);
    std::generate(data, data + n, [n = 0]() mutable { ++n; return n % 1000; });
    // Named Kernels
    oneapi::dpl::sort(oneapi::dpl::execution::make_device_policy<class sort>(q), data, data + n, [](const T& a, const T& b) { return a < b; });
    // Unnamed Kernels
    // oneapi::dpl::sort(oneapi::dpl::execution::make_device_policy(q), data, data + n, [](const T& a, const T& b) { return a < b; });

    std::cout << std::is_sorted(data, data + n) << std::endl;
    sycl::free(data, q);
}

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Projects
None yet
Development

Successfully merging this pull request may close these issues.

1 participant