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;
        });

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