Fujitsu/benchmarks/resnet/implementations/implementation_open/mxnet/src/operator/nn/cudnn/nhwc_batch_norm.h
// We calculate (based on smem usage) the achievable occupancy and make sure we run a kernel
// version that was compiled with that occupancy in its launch bounds.  This way, we avoid
// needless register spills.
We calculate (based on smem usage) the achievable occupancy and make sure we run a kernel
version that was compiled with that occupancy in its launch bounds.  This way, we avoid
needless register spills.
We calculate ( based on smem usage ) the achievable occupancy and make sure we run a kernel version that was compiled with that occupancy in its launch bounds . This way we avoid needless register spills .
void _fwdKernelLauncher(cudaStream_t stream, NhwcBatchNormFwdParams params,
                                dim3 grid_dim, int outer_loops, bool use_relu, int device_id) {
#define LAUNCH_FWD_KERNEL(OUTER_LOOPS, USE_RELU, USE_ADD_RELU, COMPILED_FOR_OCCUPANCY) \
    do { \
        CHECK(smem_size_fwd <= MAX_SMEM_WITHOUT_OPT_IN) << "Nhwc batchnorm kernel smem too big."; \
        auto fwd_func = nhwc_batch_norm_fwd< \
                        StorageType, \
                        THREADS_PER_CTA, \
                        THREADS_PER_PIXEL, \
                        PIXELS_PER_THREAD_IN_REGISTERS_FWD, \
                        ELEMENTS_PER_LDG, \
                        USE_ONLINE_APPROACH, \
                        OUTER_LOOPS, \
                        USE_RELU, \
                        USE_ADD_RELU, \
                        COMPILED_FOR_OCCUPANCY>; \
        if (COMPILED_FOR_OCCUPANCY > 1) { \
            cudaFuncSetAttribute(fwd_func, cudaFuncAttributePreferredSharedMemoryCarveout, 100); \
            checkCudaStatus(name_ + " fwd ser coop kernel (cudaFuncSetAttribute carveout)"); \
        } \
        void *params_ptr = static_cast<void*>(&params); \
        using FWD_FUNC = decltype(nhwc_batch_norm_fwd< \
                        StorageType, \
                        THREADS_PER_CTA, \
                        THREADS_PER_PIXEL, \
                        PIXELS_PER_THREAD_IN_REGISTERS_FWD, \
                        ELEMENTS_PER_LDG, \
                        USE_ONLINE_APPROACH, \
                        OUTER_LOOPS, \
                        USE_RELU, \
                        USE_ADD_RELU, \
                        COMPILED_FOR_OCCUPANCY>); \
        cudaLaunchCooperativeKernel<FWD_FUNC>(fwd_func, \
            grid_dim, \
            THREADS_PER_CTA, \
            &params_ptr, \
            smem_size_fwd, \
            stream); \
        checkCudaStatus(name_ + " fwd ser coop kernel"); \
    } while (0)
    int occupancy = smem_driven_fwd_occupancy(device_id);
    if (outer_loops == 1 && use_relu) {
      if (occupancy >= 2)
        LAUNCH_FWD_KERNEL(1, true, false, 2);
      else
        LAUNCH_FWD_KERNEL(1, true, false, 1);
    } else if (outer_loops == 1 && !use_relu) {
      if (occupancy >= 2)
        LAUNCH_FWD_KERNEL(1, false, false, 2);
      else
        LAUNCH_FWD_KERNEL(1, false, false, 1);
    } else if (use_relu) {
      if (occupancy >= 2)
        LAUNCH_FWD_KERNEL(0, true, false, 2);
      else
        LAUNCH_FWD_KERNEL(0, true, false, 1);
    } else {
      if (occupancy >= 2)
        LAUNCH_FWD_KERNEL(0, false, false, 2);
      else
        LAUNCH_FWD_KERNEL(0, false, false, 1);
    }
#undef LAUNCH_FWD_KERNEL
  }
void _fwdKernelLauncher ( cudaStream_t stream , NhwcBatchNormFwdParams params , dim3 grid_dim , int outer_loops , bool use_relu , int device_id ) { #define LAUNCH_FWD_KERNEL ( OUTER_LOOPS , USE_RELU , USE_ADD_RELU , COMPILED_FOR_OCCUPANCY )  \
    do { \
        CHECK(smem_size_fwd <= MAX_SMEM_WITHOUT_OPT_IN) << "Nhwc batchnorm kernel smem too big."; \
        auto fwd_func = nhwc_batch_norm_fwd< \
                        StorageType, \
                        THREADS_PER_CTA, \
                        THREADS_PER_PIXEL, \
                        PIXELS_PER_THREAD_IN_REGISTERS_FWD, \
                        ELEMENTS_PER_LDG, \
                        USE_ONLINE_APPROACH, \
                        OUTER_LOOPS, \
                        USE_RELU, \
                        USE_ADD_RELU, \
                        COMPILED_FOR_OCCUPANCY>; \
        if (COMPILED_FOR_OCCUPANCY > 1) { \
            cudaFuncSetAttribute(fwd_func, cudaFuncAttributePreferredSharedMemoryCarveout, 100); \
            checkCudaStatus(name_ + " fwd ser coop kernel (cudaFuncSetAttribute carveout)"); \
        } \
        void *params_ptr = static_cast<void*>(&params); \
        using FWD_FUNC = decltype(nhwc_batch_norm_fwd< \
                        StorageType, \
                        THREADS_PER_CTA, \
                        THREADS_PER_PIXEL, \
                        PIXELS_PER_THREAD_IN_REGISTERS_FWD, \
                        ELEMENTS_PER_LDG, \
                        USE_ONLINE_APPROACH, \
                        OUTER_LOOPS, \
                        USE_RELU, \
                        USE_ADD_RELU, \
                        COMPILED_FOR_OCCUPANCY>); \
        cudaLaunchCooperativeKernel<FWD_FUNC>(fwd_func, \
            grid_dim, \
            THREADS_PER_CTA, \
            &params_ptr, \
            smem_size_fwd, \
            stream); \
        checkCudaStatus(name_ + " fwd ser coop kernel"); \
    } while (0) int occupancy = smem_driven_fwd_occupancy ( device_id ) ; if ( outer_loops == 1 && use_relu ) { if ( occupancy >= 2 ) LAUNCH_FWD_KERNEL ( 1 , true , false , 2 ) ; else LAUNCH_FWD_KERNEL ( 1 , true , false , 1 ) ; } else if ( outer_loops == 1 && ! use_relu ) { if ( occupancy >= 2 ) LAUNCH_FWD_KERNEL ( 1 , false , false , 2 ) ; else LAUNCH_FWD_KERNEL ( 1 , false , false , 1 ) ; } else if ( use_relu ) { if ( occupancy >= 2 ) LAUNCH_FWD_KERNEL ( 0 , true , false , 2 ) ; else LAUNCH_FWD_KERNEL ( 0 , true , false , 1 ) ; } else { if ( occupancy >= 2 ) LAUNCH_FWD_KERNEL ( 0 , false , false , 2 ) ; else LAUNCH_FWD_KERNEL ( 0 , false , false , 1 ) ; } #undef  LAUNCH_FWD_KERNEL }
We calculate (based on smem usage) the achievable occupancy and make sure we run a kernel
version that was compiled with that occupancy in its launch bounds.
We calculate ( based on smem usage ) the achievable occupancy and make sure we run a kernel version that was compiled with that occupancy in its launch bounds .
// Don't try for an occupancy > 2 as this will squeeze register use and create spills.