Как создать 2d массив в cuda?

Question

Для практики я работаю над созданием простой программы инициализации матрицы в cuda.Я сделал небольшую последовательную версию для справки в качестве отправной точки.Он просто создает массив n на m и заполняет его значениями типа double.Я читал другие посты и документацию, но я немного сбит с толку, и я надеялся, что кто-нибудь сможет объяснить мне, как инициализировать 2d-массив в cuda способом, подобным размеру n на m, как я делал ниже.Я также был бы признателен за понимание того, как заполнить эту матрицу cuda, если кто-то захочет объяснить.

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

Последовательная версия:

#include <stdlib.h>
#include <stdio.h>
#include <math.h>
#include <assert.h>

int n,m, i, j;
double count;

void update(int n, int m, double arr[][m]){
  for(i=0; i<n; i++){
    for(j=0; j<m; j++){
      count++;
      arr[i][j] = count;
    }
  }
}


int main(int argc, char * argv[]) {
  assert(argc==3);
  n = atoi(argv[2]);
  m = atoi(argv[1]);

  double (*arr)[n][m] = malloc(sizeof *arr);
  update(n,m,arr);
  return 0;
}

Answer 1

Вы можете моделировать 2D-массив в 1D, сохраняя данные построчно.Так что двумерный массив: [a, b] [c, d] становится [a, b, c, d].Для простоты вы можете написать класс-оболочку, обеспечивающий такую ​​функциональность.

Вот демонстрационная версия (не на 100% защищенная от стихийных бедствий, но работающая) этой идеи

#pragma once
#include <iostream>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

typedef int TYPE;

// NOTE: copy consturctor and = operator need to be overloaded as well
template<class T>
struct Matrix
{
    Matrix(int r, int c) : rows(r), cols(c) {
        data = new T[r*c];
    }
    ~Matrix() {
        // As we allocated memory it needs to be freed upon destruction
        delete[] data;
        data = nullptr;
    }
    int rows, cols;
    T* data;
    T* operator[](int row) {
        // Returns pointer to "ROW", further call to [] on result will retrieve item at column in this row
        return data + (row*cols);
    }
};

// Simple cuda kernel 
__global__ void add(TYPE *a, TYPE *b, TYPE *c, int rows, int cols) {
    // Get element row and col
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    // If kernel block/grid is not sized perfectly make sure not to step outside data bounds
    if(row < rows && col < cols)
    {
        int idx = row*cols + col;
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    // m3 = m1 + m2 using cuda
    int rows = 5, cols = 5, total = rows * cols;
    Matrix<TYPE> m1{ rows,cols }, m2{ rows,cols }, m3{ rows,cols };

    // Initialization as 1D array
    for(int i = 0; i < total; i++)  {
        m1.data[i] = i;
    }

    // Or initialization as 2D array
    for(int r = 0; r < rows; r++)
        for(int c = 0; c < cols; c++)
            m2[r][c] = r*cols + c + 100;

    for(int i = 0; i < total; i++)  std::cout << m1.data[i] << ", ";
    std::cout << "\n";

    for(int r = 0; r < rows; r++) {
        for(int c = 0; c < cols; c++) 
            std::cout << m2[r][c] << ", ";
        std::cout << "\n";
    }

    // CUDA part
    TYPE *d_m1, *d_m2, *d_m3;

    // Allocation
    cudaMalloc((void **) &d_m1, total * sizeof(TYPE));
    cudaMalloc((void **) &d_m2, total * sizeof(TYPE));
    cudaMalloc((void **) &d_m3, total * sizeof(TYPE));

    // Copy m1 and m2 to GPU
    cudaMemcpy(d_m1, m1.data, total * sizeof(TYPE), cudaMemcpyHostToDevice);
    cudaMemcpy(d_m2, m2.data, total * sizeof(TYPE), cudaMemcpyHostToDevice);

    // Oversized on purpose to show row/col guard on add kernel
    dim3 grid(5, 5);
    dim3 block(5, 5);
    add <<< grid, block >>> (d_m1, d_m2, d_m3, rows, cols);

    // Copy result to m3
    cudaMemcpy(m3.data, d_m3, total * sizeof(TYPE), cudaMemcpyDeviceToHost);

    cudaFree(d_m1);
    cudaFree(d_m2);
    cudaFree(d_m3);

    for(int r = 0; r < rows; r++) {
        for(int c = 0; c < cols; c++)
            std::cout << m3[r][c] << ", ";
        std::cout << "\n";
    }

    system("pause");
    return 0;
}