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

Extension to cg scan #1522

Open
jinz2014 opened this issue Dec 13, 2023 · 2 comments
Open

Extension to cg scan #1522

jinz2014 opened this issue Dec 13, 2023 · 2 comments
Assignees
Labels
enhancement New feature or request

Comments

@jinz2014
Copy link
Contributor

jinz2014 commented Dec 13, 2023

#include <stdio.h>
#include <cooperative_groups.h>
#include <cooperative_groups/scan.h>
namespace cg = cooperative_groups;

__global__ void kernel() {
    auto thread_block = cg::this_thread_block();
    auto tile = cg::tiled_partition<8>(thread_block);
    unsigned int val = cg::inclusive_scan(tile, tile.thread_rank());
    printf("%u: %u\n", tile.thread_rank(), val);
}

Thanks

@jinz2014 jinz2014 added the enhancement New feature or request label Dec 13, 2023
@ShengchenJ
Copy link
Contributor

Hi,
We supported the tiled_partition<32>(thread_block) inclusive_scan and cooperative_scan migration. You can ref https://github.com/oneapi-src/SYCLomatic/pull/1495/files#diff-454c9fb17660120ad6cf4ce877dd77435506b17bea94bdeb8ffdf2d68d799e74

But tiled_partition<8> scan function migration has not been supported yet. You can ref the Intel experimental extension to migrate the partition group to the non-uniform group (like fixed_size_group). https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_non_uniform_groups.asciidoc

@jinz2014
Copy link
Contributor Author

Thanks.

@ShengchenJ ShengchenJ self-assigned this Dec 18, 2023
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
enhancement New feature or request
Projects
None yet
Development

No branches or pull requests

2 participants