Единая память и структура с массивами
У меня есть большой массив структур структур на CUDA, который является постоянным и доступен только для моего приложения. Очень упрощенный пример будет
struct Graph{
Node * nodes;
int nNode;
}
struct Node{
int* pos;
int nPos;
}
Моим ядрам нужно будет перемещаться по этому графику и запрашивать его. Как вы знаете, копирование этой структуры в память графического процессора с помощью cudaMalloc а также cudaMemcpy Это просто много кода, эта единая память должна устранить необходимость.
В своем коде я сгенерировал график в CPU, а затем для тестирования разработал следующее ядро
__global__ void testKernel(const Graph graph,int * d_res){
d_res[0]=graph.nNode;
};
называется как:
// using malloc for testing to make sure I know what I am doing
int * d_res,* h_res;
cudaMalloc((void **)&d_res,sizeof(int));
h_res=(int*)malloc(sizeof(int));
testKernel<<<1,1>>>(graph,d_res);
gpuErrchk( cudaPeekAtLastError() );
gpuErrchk(cudaMemcpy(h_res,d_res,sizeof(int),cudaMemcpyDeviceToHost));
с проверками ошибок отсюда.
Когда я использую testKernel как показано, он работает нормально, но если я изменю ядро на:
__global__ void testKernel(const Graph graph,int * d_res){
d_res[0]=graph.nodes[0].nPos;
};
Я получаю недопустимые ошибки доступа к памяти.
Это потому, что объединенная память неправильно обрабатывает данные этого типа? Есть ли способ убедиться, что я могу избежать записи всех явных копий в память GPU?
Полный MCVE:
#include <algorithm>
#include <cuda_runtime_api.h>
#include <cuda.h>
typedef struct node{
int* pos;
int nPos;
}Node;
typedef struct Graph{
Node * nodes;
int nNode;
}Graph;
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true)
{
if (code != cudaSuccess)
{
fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
if (abort) exit(code);
}
}
__global__ void testKernel(const Graph graph, int * d_res){
d_res[0] = graph.nNode;
// d_res[0]=graph.nodes[0].nPos; // Not working
};
int main(void){
// fake data, this comes from another process
Graph graph;
graph.nodes = (Node*)malloc(2*sizeof(Node));
graph.nNode = 2;
for (int i = 0; i < 2; i++){
// They can have different sizes in the original code
graph.nodes[i].pos = (int*)malloc(3 * sizeof(int));
graph.nodes[i].pos[0] = 0;
graph.nodes[i].pos[1] = 1;
graph.nodes[i].pos[2] = 2;
graph.nodes[i].nPos = 3;
}
printf("%d\n", graph.nNode); // Change to the kernel variable for comparison
int * d_res, *h_res;
cudaMalloc((void **)&d_res, sizeof(int));
h_res = (int*)malloc(sizeof(int));
testKernel << <1, 1 >> >(graph, d_res);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaMemcpy(h_res, d_res, sizeof(int), cudaMemcpyDeviceToHost));
printf("%d", h_res[0]);
return 0;
}
1 ответ
Ваш код не использует унифицированную память CUDA. UM не является "автоматическим" в любом случае. Это требует определенных шагов программирования, чтобы воспользоваться этим, и у него есть определенные системные требования.
Все это описано в разделе UM руководства по программированию.
Есть ли способ убедиться, что я могу избежать записи всех явных копий в память GPU?
Правильное использование единой системы обмена сообщениями должно позволить это. Вот полностью проработанный пример. Единственное, что я сделал, это механически преобразовал ваш malloc операции в хост-коде эквивалентны cudaMallocManaged операции.
$ cat t1389.cu
#include <algorithm>
#include <stdio.h>
typedef struct node{
int* pos;
int nPos;
}Node;
typedef struct Graph{
Node * nodes;
int nNode;
}Graph;
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true)
{
if (code != cudaSuccess)
{
fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
if (abort) exit(code);
}
}
__global__ void testKernel(const Graph graph, int * d_res){
d_res[0] = graph.nNode;
d_res[0]=graph.nodes[0].nPos; // Not working
};
int main(void){
// fake data, this comes from another process
Graph graph;
cudaMallocManaged(&(graph.nodes), 2*sizeof(Node));
graph.nNode = 2;
for (int i = 0; i < 2; i++){
// They can have different sizes in the original code
cudaMallocManaged(&(graph.nodes[i].pos), 3 * sizeof(int));
graph.nodes[i].pos[0] = 0;
graph.nodes[i].pos[1] = 1;
graph.nodes[i].pos[2] = 2;
graph.nodes[i].nPos = 3;
}
printf("%d\n", graph.nNode); // Change to the kernel variable for comparison
int * d_res, *h_res;
cudaMalloc((void **)&d_res, sizeof(int));
h_res = (int*)malloc(sizeof(int));
testKernel << <1, 1 >> >(graph, d_res);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaMemcpy(h_res, d_res, sizeof(int), cudaMemcpyDeviceToHost));
printf("%d", h_res[0]);
return 0;
}
$ nvcc t1389.cu -o t1389
$ cuda-memcheck ./t1389
========= CUDA-MEMCHECK
2
3========= ERROR SUMMARY: 0 errors
$
У единой системы обмена сообщениями есть ряд системных требований, которые задокументированы. Я не собираюсь читать их все здесь. В первую очередь вам нужен графический процессор cc3.0 или выше. Ваш MCVE не включал стандартную проверку ошибок, и я не пытался добавить ее. Но если у вас все еще есть проблемы с этим кодом, обязательно используйте правильную проверку ошибок CUDA и запустите ее с cuda-memcheck,
Если вся ваша структура данных, включая встроенные указатели, выделяется с помощью обычных распределителей хостов, и вы не можете это контролировать, то вы не сможете использовать ее напрямую в режиме единой системы обмена сообщениями, не выполняя какое-либо вовлеченное копирование. Исключением здесь будет система IBM Power9, как указано в разделе K.1.6 вышеупомянутого связанного раздела руководства по программированию.
Прежде чем пытаться использовать распределитель хоста (например, malloc) с UM, вы должны сначала проверить pageableMemoryAccessUsesHostPageTables собственности, как указано в этом разделе.
Это свойство в настоящее время не будет установлено ни в одной системе, кроме правильно сконфигурированной системы IBM Power9. Ни в одной системе x86 это свойство не установлено / не доступно.