code: def XNOR_gemm():
      code = r"""
  #define BLOCK_SIZE 16
  __global__ void xnor_gemm(unsigned int* A, unsigned int* B, float* C, int m, int n, int k) {
      // Block row and column
      int blockRow = blockIdx.y;
      int blockCol = blockIdx.x;
  
      // Thread row and column within Csub
      int row = threadIdx.y;
      int col = threadIdx.x;
  
      // Each thread block computes one sub-matrix Csub of C
      float* Csub = &C[BLOCK_SIZE * k * blockRow + BLOCK_SIZE * blockCol];
  
      // Shared memory used to store Asub and Bsub respectively
      __shared__ unsigned int As[BLOCK_SIZE][BLOCK_SIZE];
      __shared__ unsigned int Bs[BLOCK_SIZE][BLOCK_SIZE];
  
      // Each thread computes one element of Csub
      // by accumulating results into Cvalue
      // BLOCK_SIZE = 16 -> 256 threads, one per Csub element
      unsigned int Cvalue = 0;
  
      // Loop over all the sub-matrices of A and B that are
      // required to compute Csub
      // Multiply each pair of sub-matrices together
      // and accumulate the results
      for (int i = 0; i < (n / BLOCK_SIZE); ++i) {
  
          // Get sub-matrix Asub of A
          unsigned int* Asub = &A[BLOCK_SIZE * blockRow * n + BLOCK_SIZE * i];
  
          // Get sub-matrix Bsub of B
          unsigned int* Bsub = &B[BLOCK_SIZE * k * i + BLOCK_SIZE * blockCol];
  
          // Load Asub and Bsub from device memory to shared memory
          // Each thread loads one element of each sub-matrix
          As[row][col] = Asub[row*n+col];
          Bs[row][col] = Bsub[row*k+col];
  
          // Synchronize to make sure the sub-matrices are loaded
          // before starting the computation
          __syncthreads();
  
          // Multiply Asub and Bsub together
          for (int j = 0; j < BLOCK_SIZE; ++j) {
               Cvalue += __popc(As[row][j]^Bs[j][col]);
          }
  
          // Synchronize to make sure that the preceding
          // computation is done before loading two new
          // sub-matrices of A and B in the next iteration
          __syncthreads();
      }
  
      // Write Csub to device memory
      // Each thread writes one element
      if(col + blockCol * BLOCK_SIZE < k && row + blockRow * BLOCK_SIZE < m) {
          Csub[row * k + col] = -(2 * (float) Cvalue - 32 * n);
      }
  }
  """
  
      module = SourceModule(code)
      kernel = module.get_function("xnor_gemm")
      sig = "3P 3I"
      kernel.prepare(sig)
      return kernel
  
file_path: neon/neon/backends/kernels/cuda/binary.py
imports: from pycuda.compiler import SourceModule
  from pycuda.tools import context_dependent_memoize
label_desc: Use image node_id 1 for calling a global function with example usage: XNOR_gemm() and returns: kernel