AMReX_GpuTypes.H

Go to the documentation of this file.

#ifndef AMREX_GPU_TYPES_H_
#define AMREX_GPU_TYPES_H_
#include <AMReX_Config.H>
 
#include <AMReX_Extension.H>
 
#ifdef AMREX_USE_GPU
 
#ifdef AMREX_USE_SYCL
#  include <sycl/sycl.hpp>
#endif
 
namespace amrex {
 
#ifdef AMREX_USE_SYCL
 
struct dim3 {
    unsigned int x = 1;
    unsigned int y = 1;
    unsigned int z = 1;
    dim3 () = default;
    constexpr dim3 (unsigned int x_, unsigned int y_=1, unsigned int z_=1) : x(x_),y(y_),z(z_) {}
};
 
struct Dim1 {
    std::size_t x;
};
 
struct gpuStream_t {
    sycl::queue* queue = nullptr;
    bool operator== (gpuStream_t const& rhs) noexcept { return queue == rhs.queue; }
    bool operator!= (gpuStream_t const& rhs) noexcept { return queue != rhs.queue; }
};
 
#endif
 
}
 
#endif
 
namespace amrex::Gpu {
 
#if defined(AMREX_USE_SYCL)
 
struct Handler
{
    AMREX_GPU_HOST_DEVICE constexpr
    Handler (sycl::nd_item<1> const* a_item = nullptr, void* a_local = nullptr,
             int a_n_active_threds = -1)
        : item(a_item), local(a_local), numActiveThreads(a_n_active_threds) {}
 
    AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
    bool isFullBlock () const {
        return (numActiveThreads >= int(item->get_local_range(0)) ||
                numActiveThreads <= 0);
    }
 
    std::size_t globalIdx () const { return item->get_global_linear_id(); }
    std::size_t blockIdx () const { return item->get_group_linear_id(); }
    std::size_t threadIdx () const { return item->get_local_linear_id(); }
    //
    std::size_t gridDim () const { return item->get_group_range(0); }
    std::size_t blockDim () const { return item->get_local_range(0); }
 
    // warp index in block
    std::size_t warpIdx () const { return item->get_sub_group().get_group_id()[0]; }
    // lane index in warp
    std::size_t laneIdx () const { return item->get_sub_group().get_local_id()[0]; }
    // warp size
    std::size_t warpDim () const { return item->get_sub_group().get_group_range()[0]; }
 
    void* sharedMemory () const { return local; }
 
    void sharedBarrier () const { item->barrier(sycl::access::fence_space::local_space); }
    void globalBarrier () const { item->barrier(sycl::access::fence_space::global_space); }
    void syncThreads () const { item->barrier(sycl::access::fence_space::global_and_local); }
 
    sycl::nd_item<1> const* item;
    void* local; // SYCL shared local memory
    int numActiveThreads;
};
 
#elif defined(AMREX_USE_GPU)
 
struct Handler
{
    AMREX_GPU_HOST_DEVICE constexpr explicit Handler (int n_active_threads = -1)
        : numActiveThreads(n_active_threads) {}
 
    AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
    bool isFullBlock () const {
        return (numActiveThreads >= (int)blockDim.x ||
                numActiveThreads <= 0);
    }
 
    int numActiveThreads;
};
 
#else
 
struct Handler {};
 
#endif
 
}
 
#endif