Выход за пределы массива в cuda

Начинаю осваивать cuda и при работе с блоками и нитями задался вопросом что происходит в случае выхода за пределы массива. Тобишь, допустим, у нас есть изобаржение размером 257х257 - важно чтобы оно не было кратно числу потоков. Количество потоков тоже запишем двумерным: 16х16. Таким образом если мы укажем блок размером 16x16 (целочисленное деление, помните?), то обработается у нас не все изображение (это я даже не беру в расчет непрвильную индексацию). Таким образом нам нужно взять размер блока больше 17х17. Раз так, то встает вопрос: какого размера должен быть массив, выделяем на GPU? Так как я знаю что происходит при выходе за пределы массива на CPU, то я выделил больше памяти, нежели исходное изображение (чтобы исключить подобную ошибку). Так вот: следует ли так делать, или подобного выхода не произойдет если выделить памяти ровно столько, сколько занимает изображение? Вторую половину вопроса я бы не задавал, если бы не видел, как при попытке скопировать данные из CPU to GPU функция cudaMemcpy не выдала ошибку (invalid argument) из-за того, что я выделил массив меньше, чем собираюсь копировать (а это означает, что, в отличии от CPU, где-то сохраняется количество выделенной памяти.

Вот код (да, это альфа блендинг двух изображений):

// first is 3chanel image, second - 4 chanel image
__global__ void blend(unsigned char *first, unsigned char *second) {
  int x      = threadIdx.x + blockIdx.x * blockDim.x;
  int y      = threadIdx.y + blockIdx.y * blockDim.y;
  int width  = gridDim.x * blockDim.x;
  int offset = x + y * width;
  unsigned char *curPixelFirst  = first + offset * 3;
  unsigned char *curPixelSecond = second + offset * 4;
  if (curPixelSecond[3]) {
    float beta       = curPixelSecond[3] / 255.;
    curPixelFirst[0] = curPixelFirst[0] * (1 - beta) + curPixelSecond[0] * beta;
    curPixelFirst[1] = curPixelFirst[1] * (1 - beta) + curPixelSecond[1] * beta;
    curPixelFirst[2] = curPixelFirst[2] * (1 - beta) + curPixelSecond[2] * beta;
  }
}
void blending(cv::Mat &first, const cv::Mat &second) {
  assert(first.size == second.size);
  assert(first.elemSize() == 3);
  assert(second.elemSize() == 4);
  assert(first.elemSize1() == 1);
  assert(second.elemSize1() == 1);
  CHECK_CUDA_DEVICE();
  dim3 threads{16, 16};
  dim3 blocks{(unsigned)(first.cols + threads.x - 1) / threads.x, (unsigned)(first.rows + threads.y - 1) / threads.y};
  unsigned char *firstImageCuda  = nullptr;
  unsigned char *secondImageCuda = nullptr;
  size_t sizeOfFirst  = first.total() * first.elemSize();
  size_t sizeOfSecond = second.total() * second.elemSize();
  // here i get more memory, because result (cols / t) can have not integral part
  size_t allocSize = blocks.x * threads.x * blocks.y * threads.y;
  CUDA_THROW_IF(cudaMalloc(&firstImageCuda, allocSize * first.elemSize()));
  CUDA_THROW_IF(cudaMalloc(&secondImageCuda, allocSize * second.elemSize()));
  CUDA_THROW_IF(cudaMemcpy(firstImageCuda, first.data, sizeOfFirst, cudaMemcpyHostToDevice));
  CUDA_THROW_IF(cudaMemcpy(secondImageCuda, second.data, sizeOfSecond, cudaMemcpyHostToDevice));
  blend<<<blocks, threads>>>(firstImageCuda, secondImageCuda);
  CUDA_THROW_IF(cudaMemcpy(first.data, firstImageCuda, sizeOfFirst, cudaMemcpyDeviceToHost));
  CUDA_THROW_IF(cudaMemcpy(second.data, secondImageCuda, sizeOfSecond, cudaMemcpyDeviceToHost));
  cudaFree(firstImageCuda);
  cudaFree(secondImageCuda);
}

CUDA_THROW_IF просто бросает исключение в случае ошибки