Skip to content

Class muda::Launch

ClassList > muda > Launch

A wrapper of raw cuda kernel launch in muda style , removing the<<<>>> usage, for better intellisense support.More...

  • #include <launch.h>

Inherits the following classes: muda::LaunchBase

Public Types

Type Name
typedef KernelNodeParms< details::LaunchCallable< raw_type_t< F > > > NodeParms

Public Types inherited from muda::LaunchBase

See muda::LaunchBase

Type Name
typedef T derived_type

Public Functions

Type Name
MUDA_HOST Launch (dim3 gridDim, dim3 blockDim, size_t sharedMemSize=0, cudaStream_t stream=nullptr)
MUDA_HOST Launch (int gridDim=1, int blockDim=1, size_t sharedMemSize=0, cudaStream_t stream=nullptr)
MUDA_HOST Launch (dim3 blockDim, size_t sharedMemSize=0, cudaStream_t stream=nullptr)
MUDA_HOST Launch & apply (F && f)
MUDA_HOST Launch & apply (F && f, Tag< UserTag >)
MUDA_HOST Launch & apply (const dim3 & active_dim, F && f)
MUDA_HOST Launch & apply (const dim3 & active_dim, F && f, Tag< UserTag >)
MUDA_HOST MUDA_NODISCARD auto as_node_parms (F && f)
MUDA_HOST MUDA_NODISCARD auto as_node_parms (F && f, Tag< UserTag >)
MUDA_HOST MUDA_NODISCARD auto as_node_parms (const dim3 & active_dim, F && f)
MUDA_HOST MUDA_NODISCARD auto as_node_parms (const dim3 & active_dim, F && f, Tag< UserTag >)

Public Functions inherited from muda::LaunchBase

See muda::LaunchBase

Type Name
MUDA_GENERIC LaunchBase (::cudaStream_t stream)
T & callback (const std::function< void(::cudaStream_t, ::cudaError)> & callback)
T & file_line (std::string_view file, int line)
T & kernel_name (std::string_view name)
Next next (Next n)
Next next (Args &&... args)
T & pop_range ()
T & push_range (const std::string & name)
T & record (cudaEvent_t e, int flag=cudaEventRecordDefault)
T & record (ComputeGraphVar< cudaEvent_t > & e, const std::vector< ComputeGraphVarBase * > & vars)
T & record (ComputeGraphVar< cudaEvent_t > & e, ComputeGraphVar< ViewT > &... vars)
T & wait (cudaEvent_t e, int flag=cudaEventWaitDefault)
T & wait (const ComputeGraphVar< cudaEvent_t > & e, const std::vector< ComputeGraphVarBase * > & vars)
T & wait (const ComputeGraphVar< cudaEvent_t > & e, ComputeGraphVar< ViewT > &... vars)
T & wait ()
T & when (cudaEvent_t e, int flag=cudaEventWaitDefault)
~LaunchBase ()

Public Functions inherited from muda::LaunchCore

See muda::LaunchCore

Type Name
MUDA_GENERIC LaunchCore (::cudaStream_t stream)
void callback (const std::function< void(::cudaStream_t, ::cudaError)> & callback)
void init_stream (::cudaStream_t s)
void pop_range ()
void push_range (const std::string & name)
void record (cudaEvent_t e, int flag=cudaEventRecordDefault)
void record (ComputeGraphVar< cudaEvent_t > & e, const std::vector< ComputeGraphVarBase * > & vars)
void record (ComputeGraphVar< cudaEvent_t > & e, ComputeGraphVar< ViewT > &... vars)
void wait (cudaEvent_t e, int flag=cudaEventWaitDefault)
void wait (const ComputeGraphVar< cudaEvent_t > & e, const std::vector< ComputeGraphVarBase * > & vars)
void wait (const ComputeGraphVar< cudaEvent_t > & e, ComputeGraphVar< ViewT > &... vars)
void wait ()
void when (cudaEvent_t e, int flag=cudaEventWaitDefault)
~LaunchCore ()

Public Static Functions inherited from muda::LaunchCore

See muda::LaunchCore

Type Name
void file_line (std::string_view file, int line)
void kernel_name (std::string_view name)
void wait_device ()
void wait_event (cudaEvent_t event)
void wait_stream (::cudaStream_t stream)

Protected Types inherited from muda::LaunchCore

See muda::LaunchCore

Type Name
typedef std::shared_ptr< T > S

Protected Attributes inherited from muda::LaunchCore

See muda::LaunchCore

Type Name
::cudaStream_t m_stream

Protected Functions inherited from muda::LaunchBase

See muda::LaunchBase

Type Name
T & pop_kernel_label ()

Protected Functions inherited from muda::LaunchCore

See muda::LaunchCore

Type Name
MUDA_HOST void pop_kernel_label ()
MUDA_GENERIC::cudaStream_t stream () const

Detailed Description

A raw cuda kernel define and launch:

__global__ void cuda_kernel() {}

int main()
{
    cuda_kernel<<<4,64>>>();
}

The muda style kernel launch:

// muda kernel launch
Launch(4,64)
    .kernel_name("kernel_name") // optional
    .apply([]__device__(){}); // kernel body

A more complicated but more convincing example, to show why using muda style kernel launch is better than raw cuda kernel launch.

DeviceBuffer3D<float> volume{10,10,10};
Launch(dim3{8,8,8}) // blockDim
    .kernel_name("write_volume") // optional, for better debug info
    .apply(volume.extent(), 
        [
            volume = volume.viewer().name("volume") // name is optional, for better debug info
        ] __device__(int3 xyz) mutable
        {
            volume(xyz) = 1.0f;
        });

See also: device_buffer_3d.h parallel_for.h

Public Types Documentation

typedef NodeParms

using muda::Launch::NodeParms =  KernelNodeParms<details::LaunchCallable<raw_type_t<F> >>;

Public Functions Documentation

function Launch [1/3]

inline MUDA_HOST muda::Launch::Launch (
    dim3 gridDim,
    dim3 blockDim,
    size_t sharedMemSize=0,
    cudaStream_t stream=nullptr
) 

function Launch [2/3]

inline MUDA_HOST muda::Launch::Launch (
    int gridDim=1,
    int blockDim=1,
    size_t sharedMemSize=0,
    cudaStream_t stream=nullptr
) 

function Launch [3/3]

inline MUDA_HOST muda::Launch::Launch (
    dim3 blockDim,
    size_t sharedMemSize=0,
    cudaStream_t stream=nullptr
) 

function apply [1/4]

template<typename F, typename UserTag>
MUDA_HOST Launch & muda::Launch::apply (
    F && f
) 

function apply [2/4]

template<typename F, typename UserTag>
MUDA_HOST Launch & muda::Launch::apply (
    F && f,
    Tag < UserTag >
) 

function apply [3/4]

template<typename F, typename UserTag>
MUDA_HOST Launch & muda::Launch::apply (
    const dim3 & active_dim,
    F && f
) 

function apply [4/4]

template<typename F, typename UserTag>
MUDA_HOST Launch & muda::Launch::apply (
    const dim3 & active_dim,
    F && f,
    Tag < UserTag >
) 

function as_node_parms [1/4]

template<typename F, typename UserTag>
MUDA_HOST MUDA_NODISCARD auto muda::Launch::as_node_parms (
    F && f
) 

function as_node_parms [2/4]

template<typename F, typename UserTag>
MUDA_HOST MUDA_NODISCARD auto muda::Launch::as_node_parms (
    F && f,
    Tag < UserTag >
) 

function as_node_parms [3/4]

template<typename F, typename UserTag>
MUDA_HOST MUDA_NODISCARD auto muda::Launch::as_node_parms (
    const dim3 & active_dim,
    F && f
) 

function as_node_parms [4/4]

template<typename F, typename UserTag>
MUDA_HOST MUDA_NODISCARD auto muda::Launch::as_node_parms (
    const dim3 & active_dim,
    F && f,
    Tag < UserTag >
) 


The documentation for this class was generated from the following file src/muda/launch/launch.h