Имея тягу :: device_vector в глобальной области видимости

Question

Я пишу программу, которая рассчитывает множество свойств данных треугольной сетки.Некоторые из этих свойств, которые я хотел бы рассчитать с помощью thrust :: методов, другие свойства должны быть рассчитаны с использованием необработанных указателей памяти в ядрах CUDA.

Для передачи данных в графический процессор, я получил это в файл Transfer.cu (поскольку создание и управление thrust::device_vector s в простом коде C ++ не поддерживается):

// thrust vectors (global)
thrust::host_vector<glm::vec3> trianglethrust_host;
thrust::device_vector<glm::vec3> trianglethrust_device;

extern "C" void trianglesToGPU_thrust(const trimesh::TriMesh *mesh, float** triangles) {
// fill host vector
for (size_t i = 0; i < mesh->faces.size(); i++) {
    // PUSHING DATA INTO HOST_VECTOR HERE (OMITTED FOR CLARITY)
} 
// copy to GPU by assigning host vector to device vector, like in the Thrust documentation
trianglethrust_device = trianglethrust_host;
// save raw pointer
*triangles = (float*)thrust::raw_pointer_cast(&(trianglethrust_device[0]));
}

Эта функция trianglestoGPU_thrust вызывается из основного методамоей программы на C ++.Все работает отлично и прекрасно, пока программа не завершится, и (глобально определенный) вектор trianglethrust_device выходит из области видимости.Thrust пытается освободить его, но контекст CUDA уже исчез, в результате чего cudaErrorInvalidDevicePointer

Что будет считаться наилучшей практикой для моей проблемы?

TL; DR: Я хочу, чтобы thrust :: device_vector работал в течение всей моей программы, поскольку я хочу добавить в нее thrust :: functions (например, transform и т. Д.), А также прочитатьи манипулировать им с помощью доступа к необработанным указателям в CUDA.

Решение: В моем случае я явно освобождаюсь, используя указатель необработанных данных где-то еще в процессе.Удаление этого свободного и завершение моего основного цикла с помощью

trianglethrust_device.clear();
trianglethrust_device.shrink_to_fit();
trianglethrust_device.device_vector~;

Чтобы принудительно очистить этот вектор до того, как среда выполнения CUDA будет разрушена.Это сработало, но, вероятно, все еще довольно уродливый способ сделать это.

Я рекомендую Ответ Роберта на этот вопрос и пометит его как действительный.

Answer 1

Как вы уже обнаружили, сам контейнер вектора тяги нельзя поместить в область видимости файла.

Одно из возможных решений - просто создать нужные вам векторы в начале main, а затем передать ссылки.к этим функциям, которые им нужны.

Если вы действительно хотите «глобальное поведение», вы можете поместить указатели на векторы в глобальном / файловом контексте, затем инициализировать необходимые векторы в начале main и установить указатели наглобальная область действия, указывающая на векторы, созданные в main.

Исходя из вопроса в комментарии, я думаю, важно / желательно, чтобы основным файлом был файл .cpp, скомпилированный с компилятором хоста.Поэтому мы можем использовать ранее упомянутые концепции в сочетании с распределением векторов в куче, чтобы избежать освобождения до завершения программы.Вот полный пример:

$ cat main.cpp
#include "transfer.h"

int main(){

  float **triangles, *mesh;
  triangles = new float *[1];
  mesh = new float[4];
  mesh[0] = 0.1f; mesh[1] = 0.2f; mesh[2] = 0.3f;
  trianglesToGPU_thrust(mesh, triangles);
  do_global_work(triangles);
  finish();
}
$ cat transfer.h
void trianglesToGPU_thrust(const float *, float **);
void do_global_work(float **);
void finish();
$ cat transfer.cu
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include "transfer.h"
#include <iostream>
#include <cstdio>
#include <thrust/copy.h>

__global__ void k(float *data, size_t ds){
  for (int i = 0; i < ds; i++) printf("%f,", data[i]);
}

// thrust vectors (global)
thrust::host_vector<float> *trianglethrust_host;
thrust::device_vector<float> *trianglethrust_device;

void trianglesToGPU_thrust(const float *mesh, float** triangles) {
//create vectors
  trianglethrust_host = new thrust::host_vector<float>;
  trianglethrust_device = new thrust::device_vector<float>;

// fill host vector
  size_t i = 0;
  while (mesh[i] != 0.0f) {
    (*trianglethrust_host).push_back(mesh[i++]);
  }
// copy to GPU by assigning host vector to device vector, like in the Thrust documentation
  *trianglethrust_device = *trianglethrust_host;
// save raw pointer
  *triangles = (float*)thrust::raw_pointer_cast(&((*trianglethrust_device)[0]));
}

void do_global_work(float** triangles){

  std::cout << "from device vector:" << std::endl;
  thrust::copy((*trianglethrust_device).begin(), (*trianglethrust_device).end(), std::ostream_iterator<float>(std::cout, ","));
  std::cout << std::endl << "from kernel:" << std::endl;
  k<<<1,1>>>(*triangles, (*trianglethrust_device).size());
  cudaDeviceSynchronize();
  std::cout << std::endl;
}

void finish(){
  if (trianglethrust_host) delete trianglethrust_host;
  if (trianglethrust_device) delete trianglethrust_device;
}
$ nvcc -c transfer.cu
$ g++ -c main.cpp
$ g++ -o test main.o transfer.o -L/usr/local/cuda/lib64 -lcudart
$ ./test
from device vector:
0.1,0.2,0.3,
from kernel:
0.100000,0.200000,0.300000,
$

Вот еще один подход, аналогичный предыдущему, с использованием std::vector контейнеров тяги в глобальной области видимости (только файл transfer.cu отличается от предыдущего примера, main.cpp и transfer.h одинаковы):

$ cat transfer.cu
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include "transfer.h"
#include <iostream>
#include <cstdio>
#include <thrust/copy.h>
#include <vector>

__global__ void k(float *data, size_t ds){
  for (int i = 0; i < ds; i++) printf("%f,", data[i]);
}

// thrust vectors (global)
std::vector<thrust::host_vector<float> > trianglethrust_host;
std::vector<thrust::device_vector<float> > trianglethrust_device;

void trianglesToGPU_thrust(const float *mesh, float** triangles) {
//create vectors
  trianglethrust_host.resize(1);
  trianglethrust_device.resize(1);

// fill host vector
size_t i = 0;
  while (mesh[i] != 0.0f) {
    trianglethrust_host[0].push_back(mesh[i++]);
  }
// copy to GPU by assigning host vector to device vector, like in the Thrust documentation
  trianglethrust_device[0] = trianglethrust_host[0];
// save raw pointer
  *triangles = (float*)thrust::raw_pointer_cast(trianglethrust_device[0].data());
}

void do_global_work(float** triangles){

  std::cout << "from device vector:" << std::endl;
  thrust::copy(trianglethrust_device[0].begin(), trianglethrust_device[0].end(), std::ostream_iterator<float>(std::cout, ","));
  std::cout << std::endl << "from kernel:" << std::endl;
  k<<<1,1>>>(*triangles, trianglethrust_device[0].size());
  cudaDeviceSynchronize();
  std::cout << std::endl;
}

void finish(){
  trianglethrust_host.clear();
  trianglethrust_device.clear();
}
$ nvcc -c transfer.cu
$ g++ -o test main.o transfer.o -L/usr/local/cuda/lib64 -lcudart
$ ./test
from device vector:
0.1,0.2,0.3,
from kernel:
0.100000,0.200000,0.300000,
$