## Appendix A Implementation Variants Supported by Numba for Max Pooling Layers (Hmdl Project)
### Variant Accelerated Using GPU Computations and numba_cuda  
Another variant for max pooling computations in the hmdl project is implemented as a kernel function allowing for GPU computations. Instead of directly using the CUDA API from NVIDIA in C language, in hmdl, we take advantage of the numba_cuda module that offers analogical API and CUDA mechanisms for Python. This variant applies when the setting do_forward_impl=\"numba_cuda\" is submitted to the constructor of the MaxPool2D object. Figures A3 and A4 contain the source codes meant for that variant and forward computations. The function do_forward_numba_cuda_direct, presented in Figure A3, iterates over subbatches 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_forward_numba_cuda_direct_job, presented in Figure A4, is performed using 512 threads per block (see tpb value) and the number of blocks per grid (bpg), calculated as follows:  
\[\texttt{bpg}=\lceil m\cdot(H/S)\cdot(W/S)\cdot C/\texttt{tpb}\rceil,\]  
where \(m\) denotes the subbatch size. Figure A3: Forward computations for class MaxPool2D: variant implemented for GPU computations using numba_cuda module. ```
0$staticmethod @cuda.jit(void(float32[:,:,:,:],int32,int32,int32,int32,float32[:,:,:,:],int32[:,:,:,:])) defdo_forward_numba_cuda_direct_job(sub_input_,self_pool_size_,self_n_pools_height_,self_n_pools_width_,self_n_kernels_,sub_output_,sub_argmax_): m=sub_input_shape[@] index=cuda.grid(1) wk=self_n_pools_width_*self_n_kernels_ hwk=self_n_pools_height_*wk ifindex>=m*hwk: return i=int(index/hwk) index=int(index%hwk) p=int(index/wk) index=int(index%wk) q=int(index/self_n_kernels_) f=int(index%self_n_kernels_) p1=p*self_pool_size_ p2=p1+self_pool_size_ q1=q*self_pool_size_ q2=q1+self_pool_size_ thread_input_i_f=sub_input_[i,p1:p2,q1:q2,f] temp_max=float32(-mp.inf) temp_argmax_j=int(@) temp_argmax_k=int(@) forjinrange(self_pool_size_): forkinrange(self_pool_size_): ifthread_input_i_f[j,k]>temp_max: temp_max=thread_input_i_f[j,k] temp_argmax_j=j temp_argmax_k=k sub_output_[i,p,q,f]=temp_max sub_argmax_[i,p,q,f,theta]=p1+temp_argmax_j sub_argmax_[i,p,q,f,theta]=q1+temp_argmax_k
```
Backward GPU computations are organized analogically using the streaming mechanism (over data subbatches) and invocations of a CUDA kernel function. Figures 15 and 16 present the source codes involved in those computations. One can note how the self.argmax_information, memorized during the forward pass, is now reused at the stage of backward computations.  
defdo_backward_output_numba_cuda_direct(self):  tpb = self.cuda_max_threads_per_block_ // 2  memory = self.argmax_nbytes + self.delta_nbytes + self.delta_backward_output_.nbytes  ratio = memory / CUDA_MAX_MEMORY_PER_CALL  if ratio < 1.0:  ratio = 1.0  m_calls, call_ranges = prepare_call_ranges(self.input_shape[0],int(mp.ceil(ratio)))  streams = []  for _ _ in range(min(self.cuda_n_streams_, n_calls)):  streams.append(cuda.stream())  with cuda.pinned(self.argmax_, self.delta_, self.delta_backward_output_):  for i in range(n_calls):  stream = streams[i $ self.cuda_n_streams_]  call_slice = slice(call_ranges[i],call_ranges[i + 1])  sub_argmax = self.argmax_[call_slice]  sub_delta = self.delta_[call_slice]  dev_sub_argmax = cuda.to_device(sub_argmax, stream=stream)  dev_sub_delta = cuda.to_device(sub_delta, stream=stream)  dev_sub_delta_backward_output = cuda.device_array((dev_sub_argmax.shape[0], self. height_, self.width_, self.n_kernels_),dtype=np.float32,stream=stream)  bpg = (sub_argmax.shape[0] * self.n_pools_height_ * self.n_pools_width_ * self.