code: def _get_lut_bprop_kernel(dtype, deterministic=False):
      """
      Builds the bprop kernel for lookup table layers based on templated code.
      If the deterministic version is requested, an index buffer must be passed
      as an argument. This index buffer re-orders items in the input tensor
      so that word_ids are sorted. This is required since we need to be sure that
      each thread only updates weights for one word id.
  
      Arguments:
          dtype (np.dtype): The data which the kernel will operate on.
          deterministic (boolean): Builds the deterministic kernel when this is
              set to True.
      """
      if not deterministic:
          code = r"""
  __global__ void lut_bprop(
      int* inputs, %(type)s* dW, %(type)s* errors, const int nin,
      const int embedding_dim, const int vocab_size, const int pad_idx)
  {
      const int tid  = threadIdx.x;
      const int bid  = blockIdx.x;
  
      int word_id = inputs[bid];
      int error_row = bid * embedding_dim;
      int output_row = word_id * embedding_dim;
  
      if(word_id != pad_idx)
      {
          for(int i = tid; i < embedding_dim; i += blockDim.x)
          {
              atomicAdd(&dW[output_row + i], errors[error_row + i]);
          }
      }
  }
  """
  
          code = code % {"type": _ew_types[dtype]["type"]}
  
          module = SourceModule(code, options=["--use_fast_math"])
          kernel = module.get_function("lut_bprop")
          kernel.prepare("PPPIIIi")
      else:
          code = r"""
  __global__ void lut_bprop(
      int* inputs, int* index_buffer, %(type)s* dW, %(type)s* errors,
      const int nin, const int embedding_dim, const int vocab_size,
      const int pad_idx)
  {
      const int tid  = threadIdx.x;
      const int bid  = blockIdx.x;
  
      int index_position = bid;
      int index = index_buffer[index_position];
      int word_id = inputs[index];
  
      if((bid == 0 || word_id != inputs[index_buffer[bid - 1]]) && word_id != pad_idx)
      {
          int output_row = word_id * embedding_dim;
  
          do {
              int error_row = index * embedding_dim;
  
              for(int i = tid; i < embedding_dim; i += blockDim.x)
              {
                  dW[output_row + i] += errors[error_row + i];
              }
  
              index_position++;
              if(index_position == gridDim.x)
              {
                  break;
              }
              index = index_buffer[index_position];
          } while(inputs[index] == word_id);
      }
  }
  """
  
          code = code % {"type": _ew_types[dtype]["type"]}
  
          module = SourceModule(code, options=["--use_fast_math"])
          kernel = module.get_function("lut_bprop")
          kernel.prepare("PPPPIIIi")
  
      kernel.name = "lut_bprop"
      return kernel
  
imports: from future.utils import native_str
  from pycuda.tools import context_dependent_memoize
  from neon.backends.cuda_templates import _ew_types
  from neon.backends.util.source_module import SourceModule
label_desc: Use image node_id 1 for calling a global function with example usage: _get_lut_bprop_kernel(dtype, deterministic) and returns: kernel