Fujitsu/benchmarks/resnet/implementations/implementation_open/mxnet/src/operator/nn/cudnn/nhwc_batch_norm.h
// Helper function to launch the backward kernel.
Helper function to launch the backward kernel.
Helper function to launch the backward kernel .
void _bwdKernelLauncher(cudaStream_t stream, NhwcBatchNormBwdParams params,
                                dim3 grid_dim, int outer_loops, bool use_relu, int device_id) {
#define LAUNCH_BWD_KERNEL(OUTER_LOOPS, COMPILED_FOR_OCCUPANCY) \
    do { \
        CHECK(SMEM_SIZE_BWD <= MAX_SMEM_WITHOUT_OPT_IN) << "Nhwc batchnorm kernel smem too big."; \
        auto bwd_func = nhwc_batch_norm_bwd< \
                        StorageType, \
                        THREADS_PER_CTA, \
                        THREADS_PER_PIXEL, \
                        PIXELS_PER_THREAD_IN_REGISTERS_BWD, \
                        PIXELS_PER_THREAD_IN_SMEM_BWD, \
                        ELEMENTS_PER_LDG, \
                        USE_ONLINE_APPROACH, \
                        OUTER_LOOPS, \
                        COMPILED_FOR_OCCUPANCY>; \
        if (COMPILED_FOR_OCCUPANCY > 1) { \
            cudaFuncSetAttribute(bwd_func, cudaFuncAttributePreferredSharedMemoryCarveout, 100); \
            checkCudaStatus(name_ + " bwd coop serial kernel (cudaFuncSetAttribute carveout)"); \
        } \
        void *params_ptr = static_cast<void*>(&params); \
        using BWD_FUNC = decltype(nhwc_batch_norm_bwd< \
                        StorageType, \
                        THREADS_PER_CTA, \
                        THREADS_PER_PIXEL, \
                        PIXELS_PER_THREAD_IN_REGISTERS_BWD, \
                        PIXELS_PER_THREAD_IN_SMEM_BWD, \
                        ELEMENTS_PER_LDG, \
                        USE_ONLINE_APPROACH, \
                        OUTER_LOOPS, \
                        COMPILED_FOR_OCCUPANCY>); \
        cudaLaunchCooperativeKernel<BWD_FUNC>(bwd_func, \
            grid_dim, \
            THREADS_PER_CTA, \
            &params_ptr, \
            SMEM_SIZE_BWD, \
            stream); \
        checkCudaStatus(name_ + " bwd coop serial kernel"); \
    } while (0)
#define LAUNCH_BWD_RELU_KERNEL(OUTER_LOOPS, COMPILED_FOR_OCCUPANCY) \
    do { \
        CHECK(SMEM_SIZE_BWD <= MAX_SMEM_WITHOUT_OPT_IN) << "Nhwc batchnorm kernel smem too big."; \
        auto bwd_relu_func = nhwc_batch_norm_bwd_relu< \
                        StorageType, \
                        THREADS_PER_CTA, \
                        THREADS_PER_PIXEL, \
                        PIXELS_PER_THREAD_IN_REGISTERS_BWD, \
                        PIXELS_PER_THREAD_IN_SMEM_BWD, \
                        ELEMENTS_PER_LDG, \
                        USE_ONLINE_APPROACH, \
                        OUTER_LOOPS, \
                        COMPILED_FOR_OCCUPANCY>; \
        if (COMPILED_FOR_OCCUPANCY > 1) { \
            cudaFuncSetAttribute(bwd_relu_func, \
                cudaFuncAttributePreferredSharedMemoryCarveout, 100); \
            checkCudaStatus(name_ + \
                " bwd-relu coop serial kernel (cudaFuncSetAttribute carveout)"); \
        } \
        void *params_ptr = static_cast<void*>(&params); \
        using BWD_RELU_FUNC = decltype(nhwc_batch_norm_bwd_relu< \
                        StorageType, \
                        THREADS_PER_CTA, \
                        THREADS_PER_PIXEL, \
                        PIXELS_PER_THREAD_IN_REGISTERS_BWD, \
                        PIXELS_PER_THREAD_IN_SMEM_BWD, \
                        ELEMENTS_PER_LDG, \
                        USE_ONLINE_APPROACH, \
                        OUTER_LOOPS, \
                        COMPILED_FOR_OCCUPANCY>); \
        cudaLaunchCooperativeKernel<BWD_RELU_FUNC>(bwd_relu_func, \
            grid_dim, \
            THREADS_PER_CTA, \
            &params_ptr, \
            SMEM_SIZE_BWD, \
            stream); \
        checkCudaStatus(name_ + " bwd-relu coop serial kernel"); \
    } while (0)
    int occupancy = smem_driven_bwd_occupancy(device_id);
    if (outer_loops == 1 && use_relu) {
      if (occupancy >= 2)
        LAUNCH_BWD_RELU_KERNEL(1, 2);
      else
        LAUNCH_BWD_RELU_KERNEL(1, 1);
    } else if (outer_loops == 1 && !use_relu) {
      if (occupancy >= 2)
        LAUNCH_BWD_KERNEL(1, 2);
      else
        LAUNCH_BWD_KERNEL(1, 1);
    } else if (use_relu) {
      if (occupancy >= 2)
        LAUNCH_BWD_RELU_KERNEL(0, 2);
      else
        LAUNCH_BWD_RELU_KERNEL(0, 1);
    } else {
      if (occupancy >= 2)
        LAUNCH_BWD_KERNEL(0, 2);
      else
        LAUNCH_BWD_KERNEL(0, 1);
    }
#undef LAUNCH_BWD_KERNEL
  }
void _bwdKernelLauncher ( cudaStream_t stream , NhwcBatchNormBwdParams params , dim3 grid_dim , int outer_loops , bool use_relu , int device_id ) { #define LAUNCH_BWD_KERNEL ( OUTER_LOOPS , COMPILED_FOR_OCCUPANCY )  \
    do { \
        CHECK(SMEM_SIZE_BWD <= MAX_SMEM_WITHOUT_OPT_IN) << "Nhwc batchnorm kernel smem too big."; \
        auto bwd_func = nhwc_batch_norm_bwd< \
                        StorageType, \
                        THREADS_PER_CTA, \
                        THREADS_PER_PIXEL, \
                        PIXELS_PER_THREAD_IN_REGISTERS_BWD, \
                        PIXELS_PER_THREAD_IN_SMEM_BWD, \
                        ELEMENTS_PER_LDG, \
                        USE_ONLINE_APPROACH, \
                        OUTER_LOOPS, \
                        COMPILED_FOR_OCCUPANCY>; \
        if (COMPILED_FOR_OCCUPANCY > 1) { \
            cudaFuncSetAttribute(bwd_func, cudaFuncAttributePreferredSharedMemoryCarveout, 100); \
            checkCudaStatus(name_ + " bwd coop serial kernel (cudaFuncSetAttribute carveout)"); \
        } \
        void *params_ptr = static_cast<void*>(&params); \
        using BWD_FUNC = decltype(nhwc_batch_norm_bwd< \
                        StorageType, \
                        THREADS_PER_CTA, \
                        THREADS_PER_PIXEL, \
                        PIXELS_PER_THREAD_IN_REGISTERS_BWD, \
                        PIXELS_PER_THREAD_IN_SMEM_BWD, \
                        ELEMENTS_PER_LDG, \
                        USE_ONLINE_APPROACH, \
                        OUTER_LOOPS, \
                        COMPILED_FOR_OCCUPANCY>); \
        cudaLaunchCooperativeKernel<BWD_FUNC>(bwd_func, \
            grid_dim, \
            THREADS_PER_CTA, \
            &params_ptr, \
            SMEM_SIZE_BWD, \
            stream); \
        checkCudaStatus(name_ + " bwd coop serial kernel"); \
    } while (0) #define LAUNCH_BWD_RELU_KERNEL ( OUTER_LOOPS , COMPILED_FOR_OCCUPANCY )  \
    do { \
        CHECK(SMEM_SIZE_BWD <= MAX_SMEM_WITHOUT_OPT_IN) << "Nhwc batchnorm kernel smem too big."; \
        auto bwd_relu_func = nhwc_batch_norm_bwd_relu< \
                        StorageType, \
                        THREADS_PER_CTA, \
                        THREADS_PER_PIXEL, \
                        PIXELS_PER_THREAD_IN_REGISTERS_BWD, \
                        PIXELS_PER_THREAD_IN_SMEM_BWD, \
                        ELEMENTS_PER_LDG, \
                        USE_ONLINE_APPROACH, \
                        OUTER_LOOPS, \
                        COMPILED_FOR_OCCUPANCY>; \
        if (COMPILED_FOR_OCCUPANCY > 1) { \
            cudaFuncSetAttribute(bwd_relu_func, \
                cudaFuncAttributePreferredSharedMemoryCarveout, 100); \
            checkCudaStatus(name_ + \
                " bwd-relu coop serial kernel (cudaFuncSetAttribute carveout)"); \
        } \
        void *params_ptr = static_cast<void*>(&params); \
        using BWD_RELU_FUNC = decltype(nhwc_batch_norm_bwd_relu< \
                        StorageType, \
                        THREADS_PER_CTA, \
                        THREADS_PER_PIXEL, \
                        PIXELS_PER_THREAD_IN_REGISTERS_BWD, \
                        PIXELS_PER_THREAD_IN_SMEM_BWD, \
                        ELEMENTS_PER_LDG, \
                        USE_ONLINE_APPROACH, \
                        OUTER_LOOPS, \
                        COMPILED_FOR_OCCUPANCY>); \
        cudaLaunchCooperativeKernel<BWD_RELU_FUNC>(bwd_relu_func, \
            grid_dim, \
            THREADS_PER_CTA, \
            &params_ptr, \
            SMEM_SIZE_BWD, \
            stream); \
        checkCudaStatus(name_ + " bwd-relu coop serial kernel"); \
    } while (0) int occupancy = smem_driven_bwd_occupancy ( device_id ) ; if ( outer_loops == 1 && use_relu ) { if ( occupancy >= 2 ) LAUNCH_BWD_RELU_KERNEL ( 1 , 2 ) ; else LAUNCH_BWD_RELU_KERNEL ( 1 , 1 ) ; } else if ( outer_loops == 1 && ! use_relu ) { if ( occupancy >= 2 ) LAUNCH_BWD_KERNEL ( 1 , 2 ) ; else LAUNCH_BWD_KERNEL ( 1 , 1 ) ; } else if ( use_relu ) { if ( occupancy >= 2 ) LAUNCH_BWD_RELU_KERNEL ( 0 , 2 ) ; else LAUNCH_BWD_RELU_KERNEL ( 0 , 1 ) ; } else { if ( occupancy >= 2 ) LAUNCH_BWD_KERNEL ( 0 , 2 ) ; else LAUNCH_BWD_KERNEL ( 0 , 1 ) ; } #undef  LAUNCH_BWD_KERNEL }
Helper function to launch the backward kernel.
Helper function to launch the backward kernel .
// Don't try for an occupancy > 2 as this will squeeze register use and create spills.