Как решить библиотеку CUDA Thrust - ошибка синхронизации for_each? - PullRequest
Новая
       41

Как решить библиотеку CUDA Thrust - ошибка синхронизации for_each?

3 голосов
/ 10 марта 2012

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

Простое изменение динамического массива 1D невозможно из-за этой ошибки.

MyСегмент кода, который вызывает ошибку, выглядит следующим образом.

из файла .cpp, который я называю setIndexedGrid, который определен в System.cu 

float* a= (float*)(malloc(8*sizeof(float))); 
a[0]= 0; a[1]= 1; a[2]= 2; a[3]= 3; a[4]= 4; a[5]= 5; a[6]= 6; a[7]= 7;
float* b = (float*)(malloc(8*sizeof(float)));
setIndexedGridInfo(a,b);

Сегмент кода в System.cu:

void
setIndexedGridInfo(float* a, float*b)
{

    thrust::device_ptr<float> d_oldData(a);
    thrust::device_ptr<float> d_newData(b);

    float c = 0.0;

    thrust::for_each(
        thrust::make_zip_iterator(thrust::make_tuple(d_oldData,d_newData)),
        thrust::make_zip_iterator(thrust::make_tuple(d_oldData+8,d_newData+8)),
        grid_functor(c));
}

grid_functor определен в _kernel.cu 

struct grid_functor
{
    float a;

    __host__ __device__
    grid_functor(float grid_Info) : a(grid_Info) {}

    template <typename Tuple>
    __device__
    void operator()(Tuple t)
    {
        volatile float data = thrust::get<0>(t);
        float pos = data + 0.1;
        thrust::get<1>(t) = pos;
    }

};

Я также получаю их в окне вывода (я использую Visual Studio):

Сначалаисключение в 0x000007fefdc7cacd в Particles.exe: исключение Microsoft C ++: cudaError_enum в ячейке памяти 0x0029eb60 .. Исключение первого шанса в 0x000007fefdc7cacd в smokeParticles.exe: исключение Microsoft C ++: упор :: system :: system_error в расположении памяти 0x0029fисключение в 0x000007fefdc7cacd в Particles.exe: исключение Microsoft C ++: thrust :: system :: system_error в ячейке памяти 0x0029ecf0 ..

В чем причина проблемы?

1 Ответ

5 голосов
/ 10 марта 2012

Вы пытаетесь использовать указатели памяти хоста в функциях, ожидающих указатели в памяти устройства. Этот код является проблемой: 

float* a= (float*)(malloc(8*sizeof(float))); 
a[0]= 0; a[1]= 1; a[2]= 2; a[3]= 3; a[4]= 4; a[5]= 5; a[6]= 6; a[7]= 7;
float* b = (float*)(malloc(8*sizeof(float)));
setIndexedGridInfo(a,b);

.....

thrust::device_ptr<float> d_oldData(a);
thrust::device_ptr<float> d_newData(b);

thrust::device_ptr предназначен для "обертывания" указателя памяти устройства , выделенного с помощью API CUDA, чтобы Thrust мог его использовать. Вы пытаетесь трактовать указатель хоста напрямую как указатель устройства. Это незаконно. Вы можете изменить вашу setIndexedGridInfo функцию следующим образом: 

void setIndexedGridInfo(float* a, float*b, const int n)
{

    thrust::device_vector<float> d_oldData(a,a+n);
    thrust::device_vector<float> d_newData(b,b+n);

    float c = 0.0;

    thrust::for_each(
        thrust::make_zip_iterator(thrust::make_tuple(d_oldData.begin(),d_newData.begin())),
        thrust::make_zip_iterator(thrust::make_tuple(d_oldData.end(),d_newData.end())),
        grid_functor(c));
}

Конструктор device_vector выделит память устройства, а затем скопирует содержимое памяти вашего хоста на устройство. Это должно исправить ошибку, которую вы видите, хотя я не уверен, что вы пытаетесь сделать с помощью итератора for_each и является ли верный функтор верным.

Edit:

Вот полная, компилируемая, запускаемая версия вашего кода: 

#include <cstdlib>
#include <cstdio>
#include <thrust/device_vector.h>
#include <thrust/for_each.h>
#include <thrust/copy.h>

struct grid_functor
{
    float a;

    __host__ __device__
    grid_functor(float grid_Info) : a(grid_Info) {}

    template <typename Tuple>
    __device__
    void operator()(Tuple t)
    {
        volatile float data = thrust::get<0>(t);
        float pos = data + 0.1f;
        thrust::get<1>(t) = pos;
    }

};

void setIndexedGridInfo(float* a, float*b, const int n)
{

    thrust::device_vector<float> d_oldData(a,a+n);
    thrust::device_vector<float> d_newData(b,b+n);

    float c = 0.0;

    thrust::for_each(
        thrust::make_zip_iterator(thrust::make_tuple(d_oldData.begin(),d_newData.begin())),
        thrust::make_zip_iterator(thrust::make_tuple(d_oldData.end(),d_newData.end())),
        grid_functor(c));

    thrust::copy(d_newData.begin(), d_newData.end(), b);
}

int main(void)
{
    const int n = 8;
    float* a= (float*)(malloc(n*sizeof(float))); 
    a[0]= 0; a[1]= 1; a[2]= 2; a[3]= 3; a[4]= 4; a[5]= 5; a[6]= 6; a[7]= 7;
    float* b = (float*)(malloc(n*sizeof(float)));
    setIndexedGridInfo(a,b,n);

    for(int i=0; i<n; i++) {
        fprintf(stdout, "%d (%f,%f)\n", i, a[i], b[i]);
    }

    return 0;
}

Я могу скомпилировать и запустить этот код на хосте OS 10.6.8 с CUDA 4.1 следующим образом: 

$ nvcc -Xptxas="-v" -arch=sm_12 -g -G thrustforeach.cu 
./thrustforeach.cu(18): Warning: Cannot tell what pointer points to, assuming global memory space
./thrustforeach.cu(20): Warning: Cannot tell what pointer points to, assuming global memory space
./thrustforeach.cu(18): Warning: Cannot tell what pointer points to, assuming global memory space
./thrustforeach.cu(20): Warning: Cannot tell what pointer points to, assuming global memory space
ptxas info    : Compiling entry function '_ZN6thrust6detail7backend4cuda6detail23launch_closure_by_valueINS2_18for_each_n_closureINS_12zip_iteratorINS_5tupleINS0_15normal_iteratorINS_10device_ptrIfEEEESB_NS_9null_typeESC_SC_SC_SC_SC_SC_SC_EEEEi12grid_functorEEEEvT_' for 'sm_12'
ptxas info    : Used 14 registers, 160+0 bytes lmem, 16+16 bytes smem, 4 bytes cmem[1]
ptxas info    : Compiling entry function '_ZN6thrust6detail7backend4cuda6detail23launch_closure_by_valueINS2_18for_each_n_closureINS_12zip_iteratorINS_5tupleINS0_15normal_iteratorINS_10device_ptrIfEEEESB_NS_9null_typeESC_SC_SC_SC_SC_SC_SC_EEEEj12grid_functorEEEEvT_' for 'sm_12'
ptxas info    : Used 14 registers, 160+0 bytes lmem, 16+16 bytes smem, 4 bytes cmem[1]

$ ./a.out
0 (0.000000,0.100000)
1 (1.000000,1.100000)
2 (2.000000,2.100000)
3 (3.000000,3.100000)
4 (4.000000,4.100000)
5 (5.000000,5.100000)
6 (6.000000,6.100000)
7 (7.000000,7.100000)
...