## Appendix A Implementation Variants Supported by Numba for Max Pooling Layers (Hmdl Project)
### Variant Accelerated Using GPU Computations and numba_cuda
The outer function do_backward_numba_cuda_direct, presented in Figure A8, iterates over sub-batches within the current batch stored in self.input_, and using the CUDA streaming mechanisms, it copies the subbatches from the host to the GPU device memory. Then, an invocation of the actual CUDA kernel function, named do_backward_numba_cuda_direct_job, presented in Figure A9, is performed using 512 threads per block (see tpb value) and the number of blocks per grid (tpg), calculated as:

\[\texttt{tpg}=\lceil m\cdot F^{2}\cdot C\cdot K/\texttt{tpb}\rceil,\]

where \(m\) denotes the subbatch size. defdo_backward_numba_cuda_direct(self):  tpb=self.cuda_max_threads_per_block_ // 2  kss=self.kernel_size_*=2  total_gradient_@_batch_size =self.input_.shape[@]*kss*self.n_channels_*self.n_kernels_  memory=self.input_.mbytes+self.delta_mbytes+total_gradient_@_batch_size*np.dtype(np.float32).itemsize ratio=memory/CUDA_MAX_MEMORY_PER_CALL ifratio<1.0:  ratio=1  m_calls,call_ranges=prepare_call_ranges(self.input_.shape[@],int(np.ceil(ratio)))  streams=[]  for_inrange(min(self.cuda_n_streams_, n_calls)):  streams.append(cuda.stream())  withcuda.pinned(self.input_, self.delta_):  foriinrange(n_calls):  stream=streams[i%self.cuda_n_streams_]  call_slice=slice(call_ranges[i],call_ranges[i+1])  sub_input=self.input_[call_slice]  sub_delta=self.delta_[call_slice]  dev_sub_input=cuda.to_device(sub_input,stream=stream)  dev_sub_delta=cuda.to_device(sub_delta,stream=stream)  dsgbb_shape=(sub_input.shape[@],self.kernel_size_, self.kernel_size_, self.n_channels_, self.n_kernels_)  dev_sub_gradient_@_batch=cuda.device_array(dsgbb_shape,dtype=np.float32,  stream=stream)  bpg=(np.prod(dsg0b_shape)+tpb-1)//tpb  ConvZD.do_backward_numba_cuda_direct_job[bpg,tpb,stream](dev_sub_input,  dev_sub_delta,dev_sub_gradient_@_batch)  stream.synchronize()  self.gradient_[@]+=np.sum(dev_sub_gradient_@_batch.copy_to_host(stream=stream),  axis=@)  self.gradient_[1]=np.sum(self.delta_, axis=(@,1,2))
```  
**Figure A8**. Backward computations of gradient for class Conv2D: variant based directly on definition, using GPU computations and numba.cuda. ```
``` @staticmethod @cuda.jit(void(float32[:,:,:,:],float32[:,:,:,:],float32[:,:,:,:,:])) defdo_backward_numba_cuda_direct_job(sub_input_,sub_delta_,sub_gradient_@_batch_): m,height,width,n_channels=sub_input_shape n_kernels=sub_delta_shape[-1] kernel_size=sub_gradient_@_batch_shape[1] ksh=kernel_size>>1 index=cuda.grid(1) ck=n_channels*n_kernels ksck=kernel_size*ck ksck=kernel_size*ksck ifindex>=m*kssck: return i=int(index/kssck) index=int(index%kssck) a=int(index/ksck) a=-ksh index=int(index%ksck) b=int(index/ck) b=-ksh index=int(index%ck) l=int(index/n_kernels) s=int(index%n_kernels) temp=float32(0.0) forjinrange(height): ja=int(j+a) ifja<0 orja>=height: continue forkinrange(width): kb=int(k+b) ifkb<0 orkb>=width: continue temp+=sub_delta_[i,j,k,s]*sub_input_[i,ja,kb,1] sub_gradient_@_batch_[i,a+ksh,b+ksh,1,s]=temp
```

**Figure 9** CUDA kernel function do_backward_numba_cuda_direct_job for backward convolutional computations invoked via function do_backward_numba_cuda_direct. **Appendix B.3**: _Variant Based on Tiles Accelerated Using GPU Computations Using and numba.cuda_

The CUDA-based variant presented in this section organizes computations into _tiles_ of size \(16\times 16\). It applies 256 threads per block defined by the three-dimensional tuple \(\texttt{tpb}=(1,16,16)\) and, accordingly, the three-dimensional grid of CUDA blocks defined as follows:

\[\texttt{bpg}=(m\cdot C\cdot K,\texttt{bpg\_h},\texttt{bpg\_w})\]

where \(m\) is the subbatch size, and bpg_h and bpg_w denote the counts of tiles along the last two dimensions, both calculated as \(\lceil F/16\rceil\). Figure 10 presents the source code of the outer function (iterating over subbatches and applying the streaming CUDA mechanism).