CUDA: Выделение памяти 1d устройства для копирования массива хоста 2d указатель-указатель в и из GPU
Я работаю над проектом, пытающимся распараллелить и ускорить некоторые статистические / численные сценарии вычислений, разработанные другими людьми. До того, как начался этот проект, я был полным новичком в программировании (я больше разбираюсь в аналитической математике), поэтому, пожалуйста, прости меня за любое последующее невежество или полное недоразумение. Они используют следующую функцию для генерации матриц:
double ** CreateMatrix(int m, int n)
{
int i;
double **A;
// pointer allocation to rows
A = (double **) malloc((size_t)((m*n)*sizeof(double)));
// allocate rows and set pointers
A[0] = (double *) malloc((size_t)((m*n)*sizeof(double)));
for(i=1; i<=m; i++){
A[i]=A[i-1] + n;
}
// return the pointer to array of pointers to rows
return A;
}
Я не хочу перерабатывать базовую структуру своих матричных объектов, так как они разработали весь код вокруг него, поэтому я пытался передать эти структуры в GPU, но как 1D линейную память, как я читал распределение память и копирование указателя на массив указателей неэффективны, на GPU слишком неэффективны. Я попытался заставить этот самый простой пример работать:
__global__ void MatrixMult(double *A, double *B, double *C, int N)
{
int col = blockDim.x*blockIdx.x + threadIdx.x;
int row = blockDim.y*blockIdx.y + threadIdx.y;
if( col < N && row < N){
C[col*N + row] = A[col*N + row] + B[col*N + row];
//C[col][row] = B[col][row] + A[col][row];
}
}
const int N = 5000;
int main()
{
double **h_A,**h_B, **h_C;
h_A = CreateMatrix(N,N);
h_B = CreateMatrix(N,N);
h_C = CreateMatrix(N,N);
for(int i=0; i<N; i++){
for(int j=0; j<N; j++){
h_A[i][j]=1;
h_B[i][j]=6;
h_C[i][j]=0;
}
}
size_t pitchA,pitchB,pitchC;
double *d_A,*d_B,*d_C;
cudaMallocPitch(&d_A, &pitchA, N*sizeof(double), N);
cudaMallocPitch(&d_B, &pitchB, N*sizeof(double), N);
cudaMallocPitch(&d_C, &pitchC, N*sizeof(double), N);
cudaMemcpy2D(d_A, pitchA, h_A, N*sizeof(double), N*sizeof(double), N, cudaMemcpyHostToDevice);
cudaMemcpy2D(d_B, pitchB, h_B, N*sizeof(double), N*sizeof(double), N, cudaMemcpyHostToDevice);
cudaMemcpy2D(d_C, pitchC, h_C, N*sizeof(double), N*sizeof(double), N, cudaMemcpyHostToDevice);
dim3 GridSize(250,250,1);
dim3 BlockSize(20,20,1);
MatrixMult<<<GridSize, BlockSize>>>(d_A,d_B,d_C,N);
cudaMemcpy2D(h_C, N*sizeof(double), d_C,pitchC, N*sizeof(double), N, cudaMemcpyDeviceToHost);
PrintMatrix(h_C,N,N);
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
}
Проблема в том, что я получаю ошибку, когда пытаюсь использовать функцию PrintMatrix, чтобы проверить мои результаты:
void PrintMatrix(double **A, int m, int n)
{
int i, j;
for(i=0; i<m; i++){
for(j=0; j<n; j++){
cout << A[i][j] << "\t";
}
cout << "\n";
}
}
Я предполагаю, что есть некоторая тонкая перестройка памяти, которую я не понимаю. Я думаю, мой первый вопрос: возможно ли сдать 2D? double** объект как 1D double* на устройство, сделайте некоторые вычисления, а затем скопируйте его обратно в оригинал double** формат на хосте? Если так, может кто-нибудь сказать мне, что мне не хватает?
1 ответ
Я верю твоему CreateMatrix был сломан, но он, вероятно, был функциональным (нижеприведенная версия немного отличается от вашей, хотя ваша, вероятно, будет работать). Однако ваша общая матрица обработки между хостом и устройством была нарушена. cudaMemcpy2D а также cudaMallocPitch на самом деле не для обработки массивов двойных указателей (**) несмотря на их имена. Просмотрите документацию.
Однако ваш CreateMatrix (надлежащим образом исправлено) позволяет немного изменить ваш код и работать правильно. CreateMatrix "Умно" обеспечивает двунаправленный доступ на хосте, в то же время обеспечивая непрерывность базовых данных. Поэтому мы можем использовать A[0] в качестве указателя непосредственно на смежные данные в A, Это означает, что мы можем использовать обычные cudaMalloc а также cudaMemcpy, Вот полностью проработанный пример:
#include <iostream>
#define MAT_DIM 32
#define T1_VAL 1
#define T2_VAL 6
double ** CreateMatrix(int m, int n)
{
int i;
double **A;
// pointer allocation to rows
A = (double **) malloc((size_t)(m*sizeof(double *)));
// allocate rows and set pointers
A[0] = (double *) malloc((size_t)((m*n)*sizeof(double)));
for(i=1; i<=m; i++){
A[i]=A[i-1] + n;
}
// return the pointer to array of pointers to rows
return A;
}
void PrintMatrix(double **A, int m, int n)
{
int i, j;
for(i=0; i<m; i++){
for(j=0; j<n; j++){
std::cout << A[i][j] << "\t";
}
std::cout << "\n";
}
}
int ValidateMatrix(double **A, int m, int n)
{
int i, j;
for(i=0; i<m; i++)
for(j=0; j<n; j++)
if (A[i][j] != (T1_VAL+T2_VAL)) {printf("mismatch at %d, %d, value: %f\n", i,j,A[i][j]); return 0;}
return 1;
}
__global__ void MatrixMult(double *A, double *B, double *C, int N)
{
int col = blockDim.x*blockIdx.x + threadIdx.x;
int row = blockDim.y*blockIdx.y + threadIdx.y;
if( (col < N) && (row < N)){
C[col*N + row] = A[col*N + row] + B[col*N + row];
//C[col][row] = B[col][row] + A[col][row];
}
}
const int N = MAT_DIM;
int main()
{
double **h_A,**h_B, **h_C;
h_A = CreateMatrix(N,N);
h_B = CreateMatrix(N,N);
h_C = CreateMatrix(N,N);
for(int i=0; i<N; i++){
for(int j=0; j<N; j++){
h_A[i][j]=T1_VAL;
h_B[i][j]=T2_VAL;
h_C[i][j]=0;
}
}
double *d_A,*d_B,*d_C;
cudaMalloc(&d_A, N*N*sizeof(double));
cudaMalloc(&d_B, N*N*sizeof(double));
cudaMalloc(&d_C, N*N*sizeof(double));
cudaMemcpy(d_A, h_A[0], N*N*sizeof(double), cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B[0], N*N*sizeof(double), cudaMemcpyHostToDevice);
dim3 BlockSize(16,16);
dim3 GridSize((N+BlockSize.x-1)/BlockSize.x,(N+BlockSize.y-1)/BlockSize.y);
MatrixMult<<<GridSize, BlockSize>>>(d_A,d_B,d_C,N);
cudaMemcpy(h_C[0], d_C,N*N*sizeof(double),cudaMemcpyDeviceToHost);
//PrintMatrix(h_C,N,N);
if (!ValidateMatrix(h_C, N, N)) printf("Failure!\n");
else printf("Success!\n");
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
}
Проксимальная причина, по которой ваш PrintMatrix был segfaulting, было то, что cudaMemcpy2D операция от устройства к хосту перезаписывала массив указателей, который был установлен для индексации в h_C от CreateMatrix, Это исправлено с помощью одиночных указателей на массивы, как я показал.
Там нет ничего плохого с вашим PrintMatrix и вы должны иметь возможность раскомментировать его, если хотите. Я просто не хотел смотреть на распечатку больших матриц.
Как в стороне, ваш MatrixMult Ядро фактически добавляет 2 матрицы. Я уверен, что вы знали это.