CUDA / Проблема двойного указателя тяги (вектор указателей) - PullRequest
Новая
       4

CUDA / Проблема двойного указателя тяги (вектор указателей)

1 голос
/ 05 июня 2011

Привет всем, я использую CUDA и библиотеку Thrust.Я сталкиваюсь с проблемой, когда пытаюсь получить доступ к двойному указателю на ядре CUDA, загруженном с хоста thrust :: device_vector типа Object * (вектор указателей).При компиляции с «nvcc -o thrust main.cpp cukernel.cu» я получаю предупреждение «Предупреждение: невозможно определить, на что указывает указатель, предполагая глобальное пространство памяти» и ошибка запуска при попытке запустить программу.* Я прочитал форумы Nvidia, и похоже, что решение «Не используйте двойные указатели в ядре CUDA».Я не собираюсь сворачивать двойной указатель в одномерный указатель перед отправкой в ​​ядро ​​... Кто-нибудь нашел решение этой проблемы?Необходимый код ниже, спасибо заранее! 

--------------------------
        main.cpp
--------------------------

Sphere * parseSphere(int i)
{
  Sphere * s = new Sphere();
  s->a = 1+i;
  s->b = 2+i;
  s->c = 3+i;
  return s;
}

int main( int argc, char** argv ) {

  int i;
  thrust::host_vector<Sphere *> spheres_h;
  thrust::host_vector<Sphere> spheres_resh(NUM_OBJECTS);

  //initialize spheres_h
  for(i=0;i<NUM_OBJECTS;i++){
    Sphere * sphere = parseSphere(i);
    spheres_h.push_back(sphere);
  }

  //initialize spheres_resh
  for(i=0;i<NUM_OBJECTS;i++){
    spheres_resh[i].a = 1;
    spheres_resh[i].b = 1;
    spheres_resh[i].c = 1;
  }

  thrust::device_vector<Sphere *> spheres_dv = spheres_h;
  thrust::device_vector<Sphere> spheres_resv = spheres_resh;
  Sphere ** spheres_d = thrust::raw_pointer_cast(&spheres_dv[0]);
  Sphere * spheres_res = thrust::raw_pointer_cast(&spheres_resv[0]);

  kernelBegin(spheres_d,spheres_res,NUM_OBJECTS);

  thrust::copy(spheres_dv.begin(),spheres_dv.end(),spheres_h.begin());
  thrust::copy(spheres_resv.begin(),spheres_resv.end(),spheres_resh.begin());

  bool result = true;

  for(i=0;i<NUM_OBJECTS;i++){
    result &= (spheres_resh[i].a == i+1);
    result &= (spheres_resh[i].b == i+2);
    result &= (spheres_resh[i].c == i+3);
  }

  if(result)
  {
    cout << "Data GOOD!" << endl;
  }else{
    cout << "Data BAD!" << endl;
  }

  return 0;
}


--------------------------
        cukernel.cu
--------------------------
__global__ void deviceBegin(Sphere ** spheres_d, Sphere * spheres_res, float    
num_objects)
{
  int index = threadIdx.x + blockIdx.x*blockDim.x;

  spheres_res[index].a = (*(spheres_d+index))->a; //causes warning/launch error
  spheres_res[index].b = (*(spheres_d+index))->b; 
  spheres_res[index].c = (*(spheres_d+index))->c; 
}

void kernelBegin(Sphere ** spheres_d, Sphere * spheres_res, float num_objects)
{

 int threads = 512;//per block
 int grids = ((num_objects)/threads)+1;//blocks per grid

 deviceBegin<<<grids,threads>>>(spheres_d, spheres_res, num_objects);
}

1 Ответ

3 голосов
/ 06 июня 2011

Основная проблема в том, что вектор устройства spheres_dv содержит указатели хоста.Thrust не может выполнять «глубокое копирование» или перевод указателя между адресным пространством графического процессора и центрального процессора.Поэтому, когда вы копируете spheres_h в память GPU, вы получаете массив указателей хоста GPU.Переадресация указателей хоста на GPU недопустима - они являются указателями в неправильном адресном пространстве памяти, поэтому вы получаете эквивалентный Segfault в ядре графический процессор.

Решение будет включать в себя замену вашего parseSphere функция с чем-то, что выполняет выделение памяти на GPU, вместо использования parseSphere, который в настоящее время выделяет каждую новую структуру в памяти хоста.Если у вас был графический процессор Fermi (который, как вам кажется, у вас нет) и вы используете CUDA 3.2 или 4.0, то один из подходов - превратить parseSphere в ядро.Оператор C ++ new поддерживается в коде устройства, поэтому создание структуры будет происходить в памяти устройства.Вам нужно изменить определение Sphere, чтобы конструктор был определен как функция __device__, чтобы этот подход работал.

Альтернативный подход будет включать создание массива хоста, содержащего указатели устройства, а затем скопироватьэтот массив в память устройства.Вы можете увидеть пример этого в этом ответе .Обратите внимание, что, вероятно, случай, когда объявление thrust::device_vector, содержащее thrust::device_vector, не сработает, поэтому вам, вероятно, нужно будет создать этот массив указателей устройств, используя базовые вызовы API CUDA.

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

Суть в том, что тяга (и контейнеры C ++ STL в этом отношении) действительно не предназначены для хранения указателей.Они предназначены для хранения значений и отвлечения косвенных указателей и прямого доступа к памяти с помощью итераторов и базовых алгоритмов, которые пользователь не должен видеть.Кроме того, проблема «глубокого копирования» является основной причиной, по которой мудрые люди на форумах NVIDIA советуют использовать несколько уровней указателей в коде графического процессора.Это сильно усложняет код, а также работает медленнее на GPU.

...