divyegala

CUDA C++ and Binary Sizes

Table of Contents

1. Introduction
2. How to analyze your binary for large and/or duplicate kernels
3. CUDA Whole Compilation Mode
   • 3.1 Remove duplicate kernel instances programmatically
4. Writing shared kernel function templates in header files
   • 4.1 Anti-pattern

• 4.2 Pattern

divyegala

void cub::EmptyKernel<void>(): 8748 bytes
void raft::linalg::detail::map_kernel<1, false, float, int, raft::const_op<float>, float>(float*, int, raft::const_op<float>, float const*): 10920 bytes
void raft::linalg::detail::map_kernel<1, false, double, int, raft::const_op<double>, double>(double*, int, raft::const_op<double>, double const*): 10976 bytes
void raft::linalg::detail::map_kernel<1, false, double, int, raft::identity_op, double>(double*, int, raft::identity_op, double const*): 11304 bytes
void raft::linalg::detail::map_kernel<1, false, float, int, raft::identity_op, float>(float*, int, raft::identity_op, float const*): 11304 bytes
void raft::linalg::detail::map_kernel<1, false, float, int, ML::Solver::sgdPredictBinaryClass<float>(raft::handle_t const&, float const*, int, int, float const*, float, float*, ML::loss_funct, CUstream_st*)::{lambda(float)#2}, float>(float*, int, ML::Solver::sgdPredictBinaryClass<float>(raft::handle_t const&, float const*, int, int, float const*, float, float*, ML::loss_fun

divyegala

void raft::linalg::detail::map_kernel<1, false, float, int, raft::const_op<float>, float>(float*, int, raft::const_op<float>, float const*): 10920 bytes
void raft::linalg::detail::map_kernel<1, false, double, int, raft::const_op<double>, double>(double*, int, raft::const_op<double>, double const*): 10976 bytes
void raft::linalg::detail::map_kernel<1, false, double, int, raft::identity_op, double>(double*, int, raft::identity_op, double const*): 11304 bytes
void raft::linalg::detail::map_kernel<1, false, float, int, raft::identity_op, float>(float*, int, raft::identity_op, float const*): 11304 bytes
void raft::linalg::detail::map_kernel<1, false, float, int, raft::plug_const_op<float, raft::mul_op>, float>(float*, int, raft::plug_const_op<float, raft::mul_op>, float const*): 11424 bytes
void raft::linalg::detail::map_kernel<1, false, double, int, raft::plug_const_op<double, raft::mul_op>, double>(double*, int, raft::plug_const_op<double, raft::mul_op>, double const*): 11480 bytes
void cub::DeviceScanInitKernel

divyegala

void raft::linalg::detail::map_kernel<1, false, double, int, raft::stats::batched::detail::information_criterion<double, int>(double*, double const*, raft::stats::IC_Type, int, int, int, CUstream_st*)::{lambda(double)#1}, double>(double*, int, raft::stats::batched::detail::information_criterion<double, int>(double*, double const*, raft::stats::IC_Type, int, int, int, CUstream_st*)::{lambda(double)#1}, double const*): 11432 bytes
void cub::DeviceScanInitKernel<cub::ScanTileState<(anonymous namespace)::FillnaTemp, false> >(cub::ScanTileState<(anonymous namespace)::FillnaTemp, false>, int): 11620 bytes
void MLCommon::LinAlg::Batched::identity_matrix_kernel<double>(double*, int): 11844 bytes
void raft::linalg::detail::map_kernel<1, false, double, unsigned long, ML::_start_params(raft::handle_t&, ML::ARIMAParams<double>&, MLCommon::LinAlg::Batched::Matrix<double>&, MLCommon::LinAlg::Batched::Matrix<double> const&, ML::ARIMAOrder const&)::{lambda(double, double)#1}, double, double>(double*, unsigned long, ML::_star