/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-7.0.0/include/ck_tile/host/kernel_launch.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-7.0.0/include/ck_tile/host/kernel_launch.hpp Source File#

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-7.0.0/include/ck_tile/host/kernel_launch.hpp Source File
Go to the documentation of this file.
 // SPDX-License-Identifier: MIT
 // Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
  
 #pragma once
  
 #include "ck_tile/core/config.hpp"
 #include "ck_tile/core/utility/ignore.hpp"
 #include "ck_tile/host/hip_check_error.hpp"
 #include "ck_tile/host/stream_config.hpp"
 #include "ck_tile/host/timer.hpp"
 #include <cstddef>
 #include <hip/hip_runtime.h>
  
 namespace ck_tile {
  
 #define LOW_CU_PROCESSORS 80
 #define HIGH_CU_PROCESSORS 228
 #define OPTIMAL_LATENCY_LOW_CU_PROCESSORS 0.005
 #define OPTIMAL_LATENCY_HIGH_CU_PROCESSORS 0.0015
 #define OPTIMAL_LATENCY_SAFE_MARGIN 0.01
  
 template <int MaxThreadPerBlock, int MinBlockPerCu, typename Kernel, typename... Args>
 #if CK_TILE_USE_LAUNCH_BOUNDS
 __launch_bounds__(MaxThreadPerBlock, MinBlockPerCu)
 #endif
     __global__ void kentry(Args... args)
 {
 #if defined(__HIP_DEVICE_COMPILE__)
     Kernel{}(args...);
 #else
     (..., (ignore = args, 0));
 #endif
 }
  
 //
 // return a anonymous functor(lambda) to be called later
 // the KernelImpl should be a class without non-static data member, or let's say
 // can be instantiate with "KernelImpl{}"
 //
 // the "static __device__ operator()(some_arg)" is the entry point of KernelImpl
 //
 template <int MaxThreadPerBlock = CK_TILE_MAX_THREAD_PER_BLOCK,
           int MinBlockPerCu     = CK_TILE_MIN_BLOCK_PER_CU,
           typename KernelImpl,
           typename... Args>
 CK_TILE_HOST auto
 make_kernel(KernelImpl /*f*/, dim3 grid_dim, dim3 block_dim, std::size_t lds_byte, Args... args)
 {
     const auto kernel = kentry<MaxThreadPerBlock, MinBlockPerCu, KernelImpl, Args...>;
  
     return [=](const stream_config& s) {
         kernel<<<grid_dim, block_dim, lds_byte, s.stream_id_>>>(args...);
     };
 }
  
 template <typename... Callables>
 CK_TILE_HOST void launch_and_check(const stream_config& sc, Callables&&... callables)
 {
     // abort the sequence in case of intermediate error
     if(!((static_cast<void>(callables(sc)), hipPeekAtLastError() == hipSuccess) && ...))
     {
         HIP_CHECK_ERROR(hipGetLastError());
     }
 }
  
 // clang-format off
 /*
  * launch_kernel()
  *
  * this is the function to launch arbitrary number of kernels with optional timer(selected by stream_config)
  * the callables should have signature as "operator()(const stream_config& s){ ... }" to call
  * 
  * the simplest way is pass in a lambda function, with "[=](const stream_config& s){ call_your_kernel_here() }"
  * as signature, for the callable (pay attention to the capture list)
  * 
  * e.g.
  *  ck_tile::launch_kernel(s,
  *                      [=](const stream_config& s){ hipMemset(ptr, 0, size) },
  *                      [=](const stream_config& s){ some_kernel<<<grids, blocks>>>(arg); }
  *                      );
  * 
  * if you use ck_tile kernel, or similiar to this style (structure with "static __device__ operator()(...){}")
  * you can pass your kernel to ck_tile::make_kernel(), which will create a anonymous functor for you,
  * then pass it to ck_tile::launch_kernel()
  * 
  * e.g.
  *  ck_tile::launch_kernel(s,
  *                      ck_tile::make_kernel<T0, B0>(kernel_0{}, grids0, blocks0, 0, kargs0),
  *                      ck_tile::make_kernel<T0, B1>(kernel_1{}, grids1, blocks1, 0, kargs1),
  *                       ...);
  **/
 // clang-format on
 template <typename... Callables>
 CK_TILE_HOST float launch_kernel(const stream_config& s, Callables&&... callables)
 {
     static_assert(sizeof...(callables) > 0, "At least one callable is required!");
  
     if(!s.time_kernel_)
     {
         launch_and_check(s, std::forward<Callables>(callables)...);
         return 0;
     }
  
     auto time_launches = [&](auto timer) {
         // Warmup
         for(int i = 0; i < s.cold_niters_; i++)
         {
             launch_and_check(s, std::forward<Callables>(callables)...);
         }
  
         timer.start(s.stream_id_);
         for(int i = 0; i < s.nrepeat_; i++)
         {
             launch_and_check(s, std::forward<Callables>(callables)...);
         }
         timer.stop(s.stream_id_);
  
         return timer.duration() / s.nrepeat_;
     };
  
     if(s.is_gpu_timer_)
     {
         return time_launches(gpu_timer{});
     }
     else
     {
         return time_launches(cpu_timer{});
     }
 }
  
 template <typename PreprocessFunc, typename... Callables>
 CK_TILE_HOST float launch_kernel_preprocess(const stream_config& s,
                                             PreprocessFunc preprocess,
                                             Callables&&... callables)
 {
     static_assert(sizeof...(callables) > 0, "At least one callable is required!");
  
     if(!s.time_kernel_)
     {
         preprocess();
         launch_and_check(s, std::forward<Callables>(callables)...);
         return 0;
     }
  
     auto time_launches = [&](auto timer) {
         // Warmup
         for(int i = 0; i < s.cold_niters_; i++)
         {
             launch_and_check(s, std::forward<Callables>(callables)...);
         }
  
         timer.start(s.stream_id_);
         for(int i = 0; i < s.nrepeat_; i++)
         {
             preprocess();
             launch_and_check(s, std::forward<Callables>(callables)...);
         }
         timer.stop(s.stream_id_);
  
         hipDeviceProp_t deviceProps;
         HIP_CHECK_ERROR(hipGetDeviceProperties(&deviceProps, 0));
  
         float preprocess_offset = (deviceProps.multiProcessorCount >= HIGH_CU_PROCESSORS)
                                       ? OPTIMAL_LATENCY_HIGH_CU_PROCESSORS
                                   : (deviceProps.multiProcessorCount == LOW_CU_PROCESSORS)
                                       ? OPTIMAL_LATENCY_LOW_CU_PROCESSORS
                                       : OPTIMAL_LATENCY_SAFE_MARGIN;
         return (timer.duration() - preprocess_offset * s.nrepeat_) / s.nrepeat_;
     };
  
     if(s.is_gpu_timer_)
     {
         return time_launches(gpu_timer{});
     }
     else
     {
         return time_launches(cpu_timer{});
     }
 }
 } // namespace ck_tile