## Appendix A Implementation Variants Supported by Numba for Max Pooling Layers (Hmdl Project)
### Variant Accelerated Using GPU Computations and numba_cuda
For each subbatch, the outer function invokes the actual kernel function presented in Figure 11.

defdo_backward_numba_cuda_tiles(self):  tile_size = 16  tpb = (1, tile_size, tile_size)  bpg_h = (self.kernel_size_ + tile_size - 1) // tile_size  bpg_w = (self.kernel_size_ + tile_size - 1) // tile_size  kss = self.kernel_size_*2  total_gradient_theta_delta_size = self.input_shape[0] * kss * self.n_channels_ * self. n_kernels_  memory = self.input_.mbytes + self.delta_.mbytes + total_gradient_theta_size * np.dtype  (np.float32).timesize  ratio = memory / CUDA_MAX_MEMORY_PER_CALL  if ratio < 1.0:  ratio = 1.0  n_call_ranges = prepare_call_ranges(self.input_.shape[0], int(np.ceil(ratio)))  streams = 1  for _ in range(min(self.cuda_n_streams_, n_calls)):  streams.append(cuda.stream())  with cuda.pinned(self.weights_[0], self.weights_[1], self.input_, self.gradient_[0]):  for i in range(n_calls):  stream = streams[i% self.cuda_n_streams_]  call_slice = slice(cell_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)  dsg$b_shape = (sub_input.shape[0], self.kernel_size_, self.kernel_size_, self.  n_channels_, self.n_kernels_)  dev_sub_gradient_theta = cuda.device_array(dsg$b_shape, dtype=np.float32,  stream=stream)  bpg = (sub_input.shape[0] * self.n_channels_ * self.n_kernels_, bpg_h, bpg_w)  Conv2D.do_backward_numba_cuda_tiles_job[bpg, tpb, stream](dev_sub_input,  dev_sub_delta, dev_sub_gradient_theta)  stream.synchronize()  self.gradient_[0] += np.sum(dev_sub_gradient_theta_copy_to_host(stream=stream),  axis=0)  self.gradient_[1] = np.sum(self.delta_, axis=(0, 1, 2))```
```
0 staticmethod
```
0 cuda.jit(void(float32[:,:,:,:]),float32[:,:,:,:],float32[:,:,:,:,:])) defdo_backward_numba_cuda_tiles_job(sub_input_,sub_delta_,sub_gradient_0_batch_): shared_delta=cuda.shared.array(shape=(64,64),dtype=float32) shared_input=cuda.shared.array(shape=(80,80),dtype=float32)# padded n_kernels=sub_delta_.shape[-1] i_c_f,a,b=cuda.grid(3) ck=n_channels*n_kernels i=int(i_c_f/ck) c_f=int(i_c_f%ck) c=int(c_f/n_kernels) f=int(c_f%n_kernels) kernel_size=sub_gradient_0_batch_.shape[1] ksh=kernel_size>>1 tile_size=cuda.blockDim.y tile_size_2=tile_size*tile_size pad=tile_size ifkernel_size-1<pad: pad=kernel_size-1 height_padded=height+pad width_padded=width+pad a0=cuda.blockIdx.y*tile_size b0=cuda.blockIdx.z*tile_size ppt=int((height_padded*width_padded+tile_size_2-1)/tile_size_2) dpt=int((height*width+tile_size_2-1)/tile_size_2) dest=cuda.threadIdx.y+cuda.threadIdx.z*tile_size for_inrange(ppt): dest_j=int(dest/width_padded) dest_k=int(dest%width_padded) src_j=dest_j-ksh+a0 src_k=dest_k-ksh+b0 ifdest_j<height_padded: ifsrc_j>=0 andsrc_j<heightandsrc_k>=0 andsrc_k<width: shared_input[dest_j,dest_k]=sub_input_[i,src_j,src_k,c] else: shared_input[dest_j,dest_k]=float32(0.0) dest+=tile_size_2 dest=cuda.threadIdx.y+cuda.threadIdx.z*tile_size for_inrange(dpt): j=int(dest/width) k=int(dest%width) ifj<height: shared_delta[j,k]=sub_delta_[i,j,k,f] dest+=tile_size_2 cuda.syncthreads() temp=float32(0.0) ifa<kernel_sizeandb<kernel_size: forjinrange(height): forkinrange(width): temp+=shared_delta[j,k]*shared_input[j+a-a0,k+b-b0] ifa<kernel_sizeandb<kernel_size: sub_gradient_0_batch_[i,a,b,c,f]=temp ```

**Figure A11**. CUDA kernel function do_backward_numba_cuda_tiles_job for backward convolutional computations invoked via function do_backward_numba_cuda_tiles. ## References

* [PERSON] (1989) [PERSON] A tutorial on hidden Markov models and selected applications in speech recognition. _Proc. IEEE_**1989**, _77_, 257-286. [CrossRef]
* [PERSON] et al. (2000) [PERSON]; [PERSON]; [PERSON] Additive logistic regression: A statistical view of boosting. _Ann. Stat._**2000**, _28_, 337-407. [CrossRef]
* [PERSON] (1943) [PERSON]; [PERSON] A logical calculus of the ideas immanent in nervous activity. _Bull. Math. Biophys._**1943**, \(5\), 115-133.