 __global__ void NAME(curandStateMtgp32 *state, int size, float *result, ARG1)  \
 {                                                                              \
   int idx = blockIdx.x * BLOCK_SIZE + threadIdx.x;                             \
-  for (int i = idx; i < size; i += BLOCK_SIZE * MAX_NUM_BLOCKS) {              \
+  int rounded_size = DIVUP(size, BLOCK_SIZE) * BLOCK_SIZE;                     \
+  for (int i = idx; i < rounded_size; i += BLOCK_SIZE * MAX_NUM_BLOCKS) {      \
     float x = CURAND_FUNC(&state[blockIdx.x]);                                 \
-    x = TRANSFORM;                                                             \
-    result[i] = x;                                                             \
+    if (i < size) {                                                            \
+      x = TRANSFORM;                                                           \
+      result[i] = x;                                                           \
+    }                                                                          \
   }                                                                            \
 }
 

{
  THCudaCheck(cudaMalloc((void**)&gen->gen_states, MAX_NUM_BLOCKS * sizeof(curandStateMtgp32)));
  THCudaCheck(cudaMalloc((void**)&gen->kernel_params, sizeof(mtgp32_kernel_params)));
  if (curandMakeMTGP32Constants(mtgp32dc_params_fast_11213, gen->kernel_params) != CURAND_STATUS_SUCCESS)
  {
    THError("Creating MTGP constants failed.");
  }
}

/* Frees memory allocated during setup. */

__global__ void generate_log_normal(curandStateMtgp32 *state, int size, float *result, float mean, float stddev)
{
  int idx = blockIdx.x * BLOCK_SIZE + threadIdx.x;
  int rounded_size = DIVUP(size, BLOCK_SIZE) * BLOCK_SIZE;
  for (int i = idx; i < rounded_size; i += BLOCK_SIZE * MAX_NUM_BLOCKS) {
    float x = curand_log_normal(&state[blockIdx.x], mean, stddev);
    if (i < size) {
      result[i] = x;
    }
  }
}

