Указатель тягового устройства CUDA с аварийным завершением при копировании - PullRequest
Новая
       16

Указатель тягового устройства CUDA с аварийным завершением при копировании

0 голосов
/ 11 июня 2018

В CUDA 9.2 у меня есть что-то вроде этого: 

#ifdef __CUDA_ARCH__
    struct Context { float n[4]; } context;
#else
    typedef __m128 Context;
#endif
struct A { float k[2]; };
struct B { float q[4]; };


struct FTransform : thrust::unary_function<A, B>
{
    const Context context;


    FTransform(Context context) : context(context){}


    __device__ __host__ B operator()(const A& a) const
    {
        B b{{a.k[0], a.k[1], a.k[0]*context.n[0], a.k[1]*context.n[1]}};



        return b;
    }
};


void DoThrust(B* _bs, const Context& context, A* _as, uint32_t count)
{
    thrust::device_ptr<B> bs = thrust::device_pointer_cast(_bs);
    thrust::device_ptr<A> as = thrust::device_pointer_cast(_as);

    FTransform fTransform(context);

    auto first = thrust::make_transform_iterator(as, fTransform);
    auto last = thrust::make_transform_iterator(as + count, fTransform);


    thrust::copy(first, last, bs);
}


int main(int c, char **argv)
{
    const uint32_t Count = 4;
    Context context;

    A* as;
    B* bs;

    cudaMalloc(&as, Count*sizeof(A));
    cudaMalloc(&bs, Count*sizeof(B));

    A hostAs[Count];
    cudaMemcpy(as, hostAs, Count * sizeof(A), cudaMemcpyHostToDevice);


    DoThrust(bs, context, as, Count);


        B hostBs[Count];
        cudaMemcpy(hostBs, bs, Count * sizeof(B), cudaMemcpyDeviceToHost);//crash

return 0;
    }

Затем, когда позже я вызываю стандартный вызов cudaMemcpy (), я получаю исключение "произошел недопустимый доступ к памяти".

Если я заменяю код тяги не-эквивалентом тяги, ошибки не возникает, и все работает нормально.Различные комбинации попыток копирования в device_vectors и т. Д. Я получаю различные сбои, которые, кажется, заставляют пытаться освободить device_ptr по какой-то причине - так, может быть, это здесь по какой-то причине?

== UPDATE ==

Хорошо, это сбивает с толку, похоже, это из-за переменной-члена функтора FTransform в моем более сложном случае.Это конкретно: 

struct FTransform : thrust::unary_function<A, B>
{
    #ifdef __CUDA_ARCH__
        struct Context { float v[4]; } context;
    #else
        __m128 context;
    #endif
    ...
};

Так что я думаю, что это проблема выравнивания, как-то => на самом деле это так, как это работает: 

#ifdef __CUDA_ARCH__
    struct __align__(16) Context { float v[4]; } context;
#else
    __m128 context;
#endif

1 Ответ

0 голосов
/ 13 июня 2018

Решение должны гарантировать, что если вы используете выровненные типы в упорных элементах функторных (такие как __m128 типов SSE), которые копируются в ГПУ, что они определяются как выровнены как во время процессора и код GPU сборка NVCC в проходит - и неслучайно предположить, что даже если тип кажется естественным образом выровненным по отношению к его эквиваленту на другом проходе, он будет в порядке, так как в противном случае может произойти что-то плохое для понимания.

Так, например, _ align _ (16) необходимо в коде, подобном этому: 

struct FTransform : thrust::unary_function<A, B>
{
#ifdef __CUDA_ARCH__
    struct __align__(16) Context { float v[4]; } context;
#else
    __m128 context;
#endif


    FTransform(Context context) : context(context){}
    __device__ __host__ B operator()(const A& a) const; // function makes use of context
};
...