## Appendix A Implementation Variants Supported by Numba for Max Pooling Layers (Hmdl Project)
### Variant Accelerated Using GPU Computations and numba_cuda
n_kernels_ + tpb - 1) // tpb  MaxPool20.do_backward_output_numba_cuda_direct_job[bpg, tpb, stream](  dev_sub_argmax, dev_sub_delta, self.pool_size_, dev_sub_delta_backward_output  )  dev_sub_delta_backward_output.copy_to_host(ary-self.delta_backward_output_[  call_slice],stream=stream)  cuda.synchronize()  
Figure A5: Backward computations for class MaxPool2D: variant implemented for GPU computations using numba.cuda module. ```
```
8 staticmethod
```
8 cuda.jit(void(int32[:.;.;.;.;.;.],float32[:.;.;.;.],int32,float32[:.;.;.;.])) defdo_backward_output_numba_cuda_direct_job(sub_argmax_,sub_delta_,self_pool_size_,sub_delta_backward_output_): m,n_pools_height,n_pools_width,n_kernels,_=sub_argmax_shape index=cuda.grid(1) wk=n_pools_width*n_kernels hwk=n_pools_height*wk ifindex>=n*hwk: return i=int(index/hwk) index=int(index%hwk) p=int(index/wk) index=int(index%wk) q=int(index/n_kernels) f=int(index%n_kernels) p1=int(p*self_pool_size_) p2=int(p1+self_pool_size_) q1=int(q*self_pool_size_) q2=int(q1+self_pool_size_) forjinrange(p1,p2): forkinrange(q1,q2): sub_delta_backward_output_[i,j,k,f]=float32(0.0) amp,amq=sub_argmax_[i,p,q,f] sub_delta_backward_output_[i,amp,amq,f]=sub_delta_[i,p,q,f] ```
**Figure 6.** CUDA kernel function _do_backward_numba_cuda_direct_job for backward max pooling computations invoked via function _do_backward_numba_cuda_direct. **Appendix B. Implementation Variants Supported by Numba for Convolutional Layers (Hmdl Project)--Backward Computations**

_Appendix B.1. Variant Accelerated Using Low-Level Compilation via numba.jit_

This variant applies when the setting _do_backward_impl=numba_jit_gemm= is submitted to the constructor of the Conv2D object. Then, the function pointer self.do_backward_function_points to the _do_backward_numba_jit_gemm method presented in Figure 7. Please note the full specification of types for arguments that come after the _ejit_ decorator. defdo_backward_numba_jit_gemm(self): self.gradient_[0]=np.zeros((self.kernel_size_,self.kernel_size_,self.n_channels_,self.n_kernels_),dtype=np.float32) Conv2D_do_backward_numba_jit_job(self.input_,self.height_,self.width_,self.n_channels_,self.n_kernels_,self.kernel_size_,self.delta_,self.gradient_[0],self.gradient_[1])

8 staticmethod

9 jit(void(float32[: : : : : :], int32, int32, int32, int32, int32, int32, float32[: : : : :], float32[: : : :], float32[:], nopytbon=True, cache=True) defdo_backward_numba_jit_gemm_job(self.input_,self.height_,self.width_,self.n_channels_,self.kernel_size_,self.delta_,self.gradient_0,self.gradient_1_): ksh=self_kernel_size_, // 2 kssc=self_kernel_size_*self_n_channels_ input_padded=np.zeros((self.input_shape[0],self.height_+2*ksh,self.width_+2*ksh,self.n_channels_),dtype=np.float32) input_padded[:,ksh: -ksh,ksh: -ksh, :]=self_input_ input_image_size=self_height_*self_width_ column_tiled_inputs=np.empty((input_image_size,kssc),dtype=np.float32) input_ones=np.ones((input_image_size,1),dtype=np.float32) forirange(self.input_shape[0]): forsinrange(self.n_kernels_): sub_delta_i_s=np.copy(self.delta_[i, : : :, s]) row_delta_s=np.reshape(sub_delta_i_s, (1, input_image_size)) index=0 forainrange(-ksh,ksh+1): forbinrange(-ksh,ksh+1): forlinrange(self.n_channels_): single_input_padded_l=np.copy(input_padded[i,ksh+a:ksh+a+self.height_,ksh+b:ksh+b+self.width_,1]) column_tiled_inputs[:,index]=np.reshape(single_input_padded_l,input_image_size) index+=1 self_gradient_0[:, : :, s]+=np.reshape(row_delta_s.dot(column_tiled_inputs),(self.kernel_size_,self.kernel_size_,self.n_channels_)) self_gradient_1[s]+=row_delta_s.dot(input_ones)[0,0]

### Variant Based Directly on Definition Accelerated Using GPU Computations and numba.cuda

Similarly to the CUDA-based max pooling implementation from Appendix A.2, the implementation for convolutional layers presented here is split into two functions.