Hacktoberfest 2026: le issue che i maintainer hanno segnato per ottobre, aperte e adatte ai principianti. Sfoglia le issue Hacktoberfest

incorrect migration of __shfl_xor_sync CUDA API within a template function

Aperta
  • Qualcosa non va?
#2,189 1 commento 0 reazioni 0 assegnatari Vedi su GitHub

Nessuno ha ancora preso questa issue.

Valutazione

Difficoltà
3/5
Tempo stimato
1-2 giorni
Idoneità per principianti
45/100
Tipo di issue
Bug
Chiarezza
Abbastanza chiara
Stato di attività
Ferma
Stack tecnologico
cpp
Ambito
compilers

Direzione di ricerca

Inizia con il reproducer di template CUDA fornito ed eseguilo tramite dpct. Confronta la migrazione del riferimento KeyT basato su template con quella del riferimento esplicito int, concentrandoti sul riconoscimento di __shfl_xor_sync e sulla conseguente chiamata dpct::permute_sub_group_by_xor. Il lavoro è completato quando il caso basato su template viene migrato in modo coerente con il caso di tipo esplicito.

Scritto dal modello di indicizzazione a partire dal testo della issue.

Descrizione

bug
Describe the bug

When I tried to migrate a template function with just 1 loc in the function body.

#include <cuda_runtime.h>

template <typename KeyT, typename ValueT, uint32_t WORKGROUP_SIZE>
class WarpSort{

    __device__ static void swap(KeyT& key, ValueT& value,
                                 uint32_t const& step, uint32_t const& activemask,
                                 bool bDescending, std::true_type const& isKeyOnly)
    {
        __shfl_xor_sync(activemask, key, step, 32);

    }
}

I expect the __shfl_xor_sync would be migrated into dpct::permute_sub_group_by_xor. However, SYCLomatic does nothing to the code.
after migration:

#include <sycl/sycl.hpp>
#include <dpct/dpct.hpp>

template <typename KeyT, typename ValueT, uint32_t WORKGROUP_SIZE>
class WarpSort{

    static void swap(KeyT& key, ValueT& value,
                                 uint32_t const& step, uint32_t const& activemask,
                                 bool bDescending, std::true_type const& isKeyOnly)
    {
        __shfl_xor_sync(activemask, key, step, 32);

    }
To reproduce
#include <cuda_runtime.h>

template <typename KeyT, typename ValueT, uint32_t WORKGROUP_SIZE>
class WarpSort{

    __device__ static void swap(KeyT& key, ValueT& value,
                                 uint32_t const& step, uint32_t const& activemask,
                                 bool bDescending, std::true_type const& isKeyOnly)
    {
        __shfl_xor_sync(activemask, key, step, 32);

    }
}

run the above code with dpct

Environment
  • OS: Linux
  • Target device and vendor: Nvidia GPU
  • DPC++ version:Intel(R) oneAPI DPC++/C++ Compiler 2024.2.0 (2024.2.0.20240602)
Additional context

One interesting observation is when you change the type of key into explicit type name such as int& key, the migration success.
before migration:

#include <cuda_runtime.h>

template <typename KeyT, typename ValueT, uint32_t WORKGROUP_SIZE>
class WarpSort{

    __device__ static void swap(int& key, ValueT& value,
                                 uint32_t const& step, uint32_t const& activemask,
                                 bool bDescending, std::true_type const& isKeyOnly)
    {
        __shfl_xor_sync(activemask, key, step, 32);

    }
}

after migration:

template <typename KeyT, typename ValueT, uint32_t WORKGROUP_SIZE>
class WarpSort{

    static void swap(int& key, ValueT& value,
                                 uint32_t const& step, uint32_t const& activemask,
                                 bool bDescending, std::true_type const& isKeyOnly,
                                 const sycl::nd_item<3> &item_ct1)
    {
        /*
        DPCT1023:0: The SYCL sub-group does not support mask options for
        dpct::permute_sub_group_by_xor. You can specify
        "--use-experimental-features=masked-sub-group-operation" to use the
        experimental helper function to migrate __shfl_xor_sync.
        */
        dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), key, step);
    }
Lingua principale
LLVM
Stelle
292
Fork
99
Merge medio
5g 5h
PR unite (30g)
1

Preparare l'ambiente

Come iniziare

  1. Leggi tutta la issue e poi la guida ai contributi del progetto.
  2. Commenta sulla issue per dire che te ne occupi tu — evita che due persone facciano lo stesso lavoro.
  3. Fai un fork del repository e lavora su un branch.
  4. Apri una pull request che faccia riferimento al numero della issue.

Altre issue di oneapi-src/SYCLomatic

Tutte le issue di oneapi-src/SYCLomatic

Issue simili

Altre issue su Compilers

Ricevi le nuove issue nella tua casella

Un breve riepilogo di issue GitHub adatte ai principianti.