text: <%
  if is_int && !is_object
    if /Int64$/ =~ class_name
      rand_bit = 64
    else
      rand_bit = 32
    end
    m_rand = "m_rand(max,shift)"
    shift_def = "int shift;"
    shift_set = "shift = #{rand_bit-1} - msb_pos(max);"
    rand_type = "uint#{rand_bit}_t"
  %>
  
  #define HWID (sizeof(dtype)*4)
  
  static int msb_pos(<%=rand_type%> a)
  {
      int width = HWID;
      int pos = 0;
      <%=rand_type%> mask = (((dtype)1 << HWID)-1) << HWID;
  
      if (a==0) {return -1;}
  
      while (width) {
          if (a & mask) {
              pos += width;
          } else {
              mask >>= width;
          }
          width >>= 1;
          mask &= mask << width;
      }
      return pos;
  }
  
  /* generates a random number on [0,max) */
  <% if rand_bit == 64 %>
  inline static dtype m_rand(uint64_t max, int shift)
  {
      uint64_t x;
      do {
          x = gen_rand32();
          x <<= 32;
          x |= gen_rand32();
          x >>= shift;
      } while (x >= max);
      return x;
  }
  <% else %>
  inline static dtype m_rand(uint32_t max, int shift)
  {
      uint32_t x;
      do {
          x = gen_rand32();
          x >>= shift;
      } while (x >= max);
      return x;
  }
  <% end %>
  <%
  else
    m_rand = "m_rand(max)"
    shift_def = ""
    shift_set = ""
    rand_type = "dtype"
  end
  %>
  
  typedef struct {
      dtype low;
      <%=rand_type%> max;
  } rand_opt_t;
  
  static void
  <%=c_iter%>(cumo_na_loop_t *const lp)
  {
      size_t   i;
      char    *p1;
      ssize_t  s1;
      size_t  *idx1;
      dtype    x;
      rand_opt_t *g;
      dtype    low;
      <%=rand_type%> max;
      <%=shift_def%>
  
      CUMO_INIT_COUNTER(lp, i);
      CUMO_INIT_PTR_IDX(lp, 0, p1, s1, idx1);
      g = (rand_opt_t*)(lp->opt_ptr);
      low = g->low;
      max = g->max;
      <%=shift_set%>
  
      CUMO_SHOW_SYNCHRONIZE_FIXME_WARNING_ONCE("<%=name%>", "<%=type_name%>");
      cumo_cuda_runtime_check_status(cudaDeviceSynchronize());
      if (idx1) {
          for (; i--;) {
              x = m_add(<%=m_rand%>,low);
              CUMO_SET_DATA_INDEX(p1,idx1,dtype,x);
          }
      } else {
          for (; i--;) {
              x = m_add(<%=m_rand%>,low);
              CUMO_SET_DATA_STRIDE(p1,s1,dtype,x);
          }
      }
  }
  
  
  /*
    Generate uniformly distributed random numbers on self narray.
    @overload rand([[low],high])
    @param [Numeric] low  lower inclusive boundary of random numbers. (default=0)
    @param [Numeric] high  upper exclusive boundary of random numbers. (default=1 or 1+1i for complex types)
    @return [Cumo::<%=class_name%>] self.
    @example
      Cumo::DFloat.new(6).rand
      => Cumo::DFloat#shape=[6]
         [0.0617545, 0.373067, 0.794815, 0.201042, 0.116041, 0.344032]
      Cumo::DComplex.new(6).rand(5+5i)
      => Cumo::DComplex#shape=[6]
         [2.69974+3.68908i, 0.825443+0.254414i, 0.540323+0.34354i, 4.52061+2.39322i, ...]
      Cumo::Int32.new(6).rand(2,5)
      => Cumo::Int32#shape=[6]
         [4, 3, 3, 2, 4, 2]
  */
  static VALUE
  <%=c_func(-1)%>(int argc, VALUE *args, VALUE self)
  {
      rand_opt_t g;
      VALUE v1=Qnil, v2=Qnil;
      dtype high;
      cumo_ndfunc_arg_in_t ain[1] = {{CUMO_OVERWRITE,0}};
      cumo_ndfunc_t ndf = {<%=c_iter%>, CUMO_FULL_LOOP, 1,0, ain,0};
  
      <% if is_int && !is_object %>
      rb_scan_args(argc, args, "11", &v1, &v2);
      if (v2==Qnil) {
          g.low = m_zero;
          g.max = high = m_num_to_data(v1);
      <% else %>
      rb_scan_args(argc, args, "02", &v1, &v2);
      if (v2==Qnil) {
          g.low = m_zero;
          if (v1==Qnil) {
              <% if is_complex %>
              g.max = high = c_new(1,1);
              <% else %>
              g.max = high = m_one;
              <% end %>
          } else {
              g.max = high = m_num_to_data(v1);
          }
      <% end %>
      } else {
          g.low = m_num_to_data(v1);
          high = m_num_to_data(v2);
          g.max = m_sub(high,g.low);
      }
      <% if is_int && !is_object %>
      if (high <= g.low) {
          rb_raise(rb_eArgError,"high must be larger than low");
      }
      <% end %>
      cumo_na_ndloop3(&ndf, &g, 1, self);
      return self;
  }
  