Выделить 2D массив на память устройства в CUDA

Как выделить и перенести (на хост и с хоста) 2D-массивы в память устройства в Cuda?

Источник

4 ответа

Решение

Я нашел решение этой проблемы. Мне не нужно было выравнивать массив.

Встроенный cudaMallocPitch() Функция сделала работу. И я мог бы передать массив и с устройства, используя cudaMemcpy2D() функция.

Например

cudaMallocPitch((void**) &array, &pitch, a*sizeof(float), b);

Это создает двумерный массив размером a * b с шагом, передаваемым в качестве параметра.

Следующий код создает двумерный массив и перебирает элементы. Он легко компилируется, вы можете использовать его.

#include<stdio.h>
#include<cuda.h>
#define height 50
#define width 50

// Device code
__global__ void kernel(float* devPtr, int pitch)
{
    for (int r = 0; r < height; ++r) {
        float* row = (float*)((char*)devPtr + r * pitch);
        for (int c = 0; c < width; ++c) {
             float element = row[c];
        }
    }
}

//Host Code
int main()
{

float* devPtr;
size_t pitch;
cudaMallocPitch((void**)&devPtr, &pitch, width * sizeof(float), height);
kernel<<<100, 512>>>(devPtr, pitch);
return 0;
}
Источник

Свести это: сделать его одномерным. Посмотрите, как это сделано здесь

Источник

Код вашего устройства может быть быстрее. Попробуйте использовать темы больше.

__global__ void kernel(float* devPtr, int pitch)
{
    int r = threadIdx.x;

    float* row = (float*)((char*)devPtr + r * pitch);
    for (int c = 0; c < width; ++c) {
         float element = row[c];
    }
}

Затем вы рассчитываете распределение блоков и потоков так, чтобы каждый поток имел дело с одним элементом.

Источник 
      #include <stdio.h>
#include <stdlib.h>
#include <sys/time.h>
#include <cuda.h>
#define MAX_ITER 1000000
#define MAX 100 //maximum value of the matrix element
#define TOL 0.000001

// Generate a random float number with the maximum value of max
float rand_float(int max){
  return ((float)rand()/(float)(RAND_MAX)) * max;
}

__global__ void kernel(float **device_2Darray1, float **device_2Darray2, float **device_2Darray3, int rows, int cols) {

  // Calculate the row index
  int row = blockIdx.y * blockDim.y + threadIdx.y;

  // Calculate the column index
  int col = blockIdx.x * blockDim.x + threadIdx.x;
  
  // Check if the thread is within the array bounds
  if (row < rows && col < cols) {
    // Perform the computation
    device_2Darray3[row][col] = device_2Darray1[row][col] + device_2Darray2[row][col];
  }
}


int main(int argc, char *argv[]){
  float **host_2Darray;
  float **device_2Darray;

  int rows = 10; // or whatever value you want
  int cols = 10; // or whatever value you want

  // allocate memory for the host
  host_2Darray = (float**)malloc(rows * sizeof(float*));
  for(int i = 0; i < rows; i++){
    host_2Darray[i] = (float*)malloc(cols * sizeof(float));
    for(int j = 0; j < cols; j++){
      host_2Darray[i][j] = rand_float(MAX);
    }
  }

  // allocate memory for the device
  cudaMalloc((void***)&device_2Darray, rows * sizeof(float*));
  for(int i = 0; i < rows; i++){
    cudaMalloc((void**)&device_2Darray[i], cols * sizeof(float));
  }

  // copy host memory to device
  for(int i = 0; i < rows; i++){
    cudaMemcpy(device_2Darray[i], host_2Darray[i], cols * sizeof(float), cudaMemcpyHostToDevice);
  }

  // call the kernel
  dim3 threadsPerBlock(16, 16);
  dim3 blocksPerGrid((rows + threadsPerBlock.x - 1) / threadsPerBlock.x, 
                     (cols + threadsPerBlock.y - 1) / threadsPerBlock.y);
  kernel<<<blocksPerGrid, threadsPerBlock>>>(device_2Darray, rows, cols);

  // copy device memory back to host
  for(int i = 0; i < rows; i++){
    cudaMemcpy(host_2Darray[i], device_2Darray[i], cols * sizeof(float), cudaMemcpyDeviceToHost);
  }

  // free device memory
  for(int i = 0; i < rows; i++){
    cudaFree(device_2Darray[i]);
  }
  cudaFree(device_2Darray);

  // free host memory
  for(int i = 0; i < rows; i++){
    free(host_2Darray[i]);
  }
  free(host_2Darray);

  return 0;
}

Метод создания 2D-массивов в CUDA сложнее, чем у 1D-массива, поскольку память устройства (GPU) линейна. Фактически мы создаем массив указателей (каждый указатель указывает на одномерный массив), отсюда и двойные указатели.

Что касается глобальной функции ядра void: простая операция, такая как поэлементное добавление двух 2D-массивов. Вы можете использовать BlockIdx и BlockDim вместе с ThreadIdx для вычисления глобального индекса потока, а затем использовать этот индекс, чтобы указать, за какой элемент массива отвечает этот поток. В этом примере BlockIdx.y и BlockIdx.x дают индекс текущего потока внутри блока в направлениях y и x. Вы можете представить это как трехуровневую иерархию: сетка -> блок -> поток.

Вам нужно будет настроить остальную часть вашего кода, чтобы выделить и инициализировать второй массив, выделить выходной массив, скопировать выходные данные обратно на хост и т. д.

Не забудьте также изменить вызов ядра, чтобы передать правильные аргументы. Предполагается, что часть работы вы проделаете самостоятельно, например, интеграцию концепций, поскольку никто не имеет доступа к вашему полному коду и тому, как все это работает вместе как система.

Источник
Другие вопросы по тегам