class_docstr: .. versionadded:: 2016.2
  
  .. automethod:: fill_uniform
  .. automethod:: uniform
  .. automethod:: fill_normal
  .. automethod:: normal
class_docstr_tok: ["..","versionadded",":",":","2016.2","..","automethod",":",":","fill_uniform","..","automethod",":",":","uniform","..","automethod",":",":","fill_normal","..","automethod",":",":","normal"]
code: def get_gen_kernel(self, dtype, distribution):
      size_multiplier = 1
      arg_dtype = dtype
  
      rng_key = (distribution, dtype)
  
      if rng_key in [("uniform", np.float64), ("normal", np.float64)]:
          c_type = "double"
          scale1_const = "((double) %r)" % (1 / 2**32)
          scale2_const = "((double) %r)" % (1 / 2**64)
          if distribution == "normal":
              transform = "box_muller"
          else:
              transform = ""
  
          rng_expr = (
              "shift + scale * "
              "%s( %s * convert_double4(gen)"
              "+ %s * convert_double4(gen))"
              % (transform, scale1_const, scale2_const)
          )
  
          counter_multiplier = 2
  
      elif rng_key in [
          (dist, cmp_dtype)
          for dist in ["normal", "uniform"]
          for cmp_dtype in [
              np.float32,
              cltypes.float2,
              cltypes.float3,
              cltypes.float4,
          ]
      ]:
          c_type = "float"
          scale_const = "((float) %r)" % (1 / 2**32)
  
          if distribution == "normal":
              transform = "box_muller"
          else:
              transform = ""
  
          rng_expr = "shift + scale * %s(%s * convert_float4(gen))" % (
              transform,
              scale_const,
          )
          counter_multiplier = 1
          arg_dtype = np.float32
          try:
              _, size_multiplier = cltypes.vec_type_to_scalar_and_count[
                  dtype
              ]
          except KeyError:
              pass
  
      elif rng_key == ("uniform", np.int32):
          c_type = "int"
          rng_expr = (
              "shift + convert_int4((convert_long4(gen) * scale) / %s)"
              % (str(2**32) + "l")
          )
          counter_multiplier = 1
  
      elif rng_key == ("uniform", np.int64):
          c_type = "long"
          rng_expr = (
              "shift"
              "+ convert_long4(gen) * (scale/two32) "
              "+ ((convert_long4(gen) * scale) / two32)".replace(
                  "two32", (str(2**32) + "l")
              )
          )
          counter_multiplier = 2
  
      else:
          raise TypeError(
              "unsupported RNG distribution/data type combination '%s/%s'"
              % rng_key
          )
  
      kernel_name = f"rng_gen_{self.generator_name}_{distribution}"
      src = """//CL//
              #include <{header_name}>
  
              #ifndef M_PI
              #ifdef M_PI_F
              #define M_PI M_PI_F
              #else
              #define M_PI 3.14159265359f
              #endif
              #endif
  
              typedef {output_t} output_t;
              typedef {output_t}4 output_vec_t;
              typedef {gen_name}_ctr_t ctr_t;
              typedef {gen_name}_key_t key_t;
  
              uint4 gen_bits(key_t *key, ctr_t *ctr)
              {{
                  union {{
                      ctr_t ctr_el;
                      uint4 vec_el;
                  }} u;
  
                  u.ctr_el = {gen_name}(*ctr, *key);
                  if (++ctr->v[0] == 0)
                      if (++ctr->v[1] == 0)
                          ++ctr->v[2];
  
                  return u.vec_el;
              }}
  
              #if {include_box_muller}
              output_vec_t box_muller(output_vec_t x)
              {{
                  #define BOX_MULLER(I, COMPA, COMPB) \
                      output_t r##I = sqrt(-2*log(x.COMPA)); \
                      output_t c##I; \
                      output_t s##I = sincos((output_t) (2*M_PI) * x.COMPB, &c##I);
  
                  BOX_MULLER(0, x, y);
                  BOX_MULLER(1, z, w);
                  return (output_vec_t) (r0*c0, r0*s0, r1*c1, r1*s1);
              }}
              #endif
  
              #define GET_RANDOM_NUM(gen) {rng_expr}
  
              kernel void {kernel_name}(
                  int k1,
                  #if {key_length} > 2
                  int k2, int k3,
                  #endif
                  int c0, int c1, int c2, int c3,
                  global output_t *output,
                  long out_size,
                  output_t scale,
                  output_t shift)
              {{
                  #if {key_length} == 2
                  key_t k = {{{{get_global_id(0), k1}}}};
                  #else
                  key_t k = {{{{get_global_id(0), k1, k2, k3}}}};
                  #endif
  
                  ctr_t c = {{{{c0, c1, c2, c3}}}};
  
                  // output bulk
                  unsigned long idx = get_global_id(0)*4;
                  while (idx + 4 < out_size)
                  {{
                      output_vec_t ran = GET_RANDOM_NUM(gen_bits(&k, &c));
                      vstore4(ran, 0, &output[idx]);
                      idx += 4*get_global_size(0);
                  }}
  
                  // output tail
                  output_vec_t tail_ran = GET_RANDOM_NUM(gen_bits(&k, &c));
                  if (idx < out_size)
                    output[idx] = tail_ran.x;
                  if (idx+1 < out_size)
                    output[idx+1] = tail_ran.y;
                  if (idx+2 < out_size)
                    output[idx+2] = tail_ran.z;
                  if (idx+3 < out_size)
                    output[idx+3] = tail_ran.w;
              }}
              """.format(
          kernel_name=kernel_name,
          gen_name=self.generator_name,
          header_name=self.header_name,
          output_t=c_type,
          key_length=self.key_length,
          include_box_muller=int(distribution == "normal"),
          rng_expr=rng_expr,
      )
  
      prg = cl.Program(self.context, src).build()
      knl = getattr(prg, kernel_name)
      knl.set_scalar_arg_dtypes(
          [np.int32] * (self.key_length - 1 + 4)
          + [None, np.int64, arg_dtype, arg_dtype]
      )
  
      return knl, counter_multiplier, size_multiplier
  
imports: import pyopencl
  import pyopencl.array
  import pyopencl.cltypes
  from pyopencl.tools import first_arg_dependent_memoize
  from pytools import memoize_method
  import numpy
label_desc: Use image node_id 5 for calling the Random123GeneratorBase obj's underlying member method code with example usage: obj.get_gen_kernel(dtype, distribution) and returns: knl,  counter_multiplier,  size_multiplier