cuda nvcc сделать __device__ условным - PullRequest
Новая
       8

cuda nvcc сделать __device__ условным

0 голосов
/ 03 сентября 2018

Я пытаюсь добавить серверную часть cuda в библиотеку шаблонов выражений c ++ с 20k loc. Пока он работает отлично, но я полностью утончен в «предупреждении: вызов функции __host__ из функции __host__ __device__ не разрешен».

Большая часть кода может быть обобщена следующим образом: 

template<class Impl>
struct Wrapper{
    Impl impl;
    // lots and lots of decorator code
    __host__ __device__ void call(){ impl.call();};
};


//Guaranteed to never ever be used on gpu.
struct ImplCPU{
    void call();
};
//Guaranteed to never ever be used on cpu.
struct ImplGPU{
    __host__ __device__ void call();//Actually only __device__, but needed to shut up the compiler as well
};

Wrapper<ImplCPU> wrapCPU;
Wrapper<ImplGPU> wrapGPU;

Во всех случаях call () в Wrapper тривиален, а сама оболочка - довольно сложный зверь (только хост-функции, содержащие метаинформацию). условная компиляция не является опцией, оба пути предназначены для использования рядом.

Мне не хватает «--disable-warnings», потому что, честно говоря, затраты на копирование и поддержку 10 тыс. Лок ужасной магии шаблонов перевешивают преимущества предупреждений.

Я был бы очень рад, если бы вызов был device или host , условно зависящий от того, предназначена ли реализация для gpu или cpu (потому что Impl знает, для чего он)

Просто чтобы показать, что это плохо. Одно предупреждение: 

/home/user/Remora/include/remora/detail/matrix_expression_classes.hpp(859): warning: calling a __host__ function from a __host__ __device__ function is not allowed
          detected during:
            instantiation of "remora::matrix_matrix_prod<MatA, MatB>::size_type remora::matrix_matrix_prod<MatA, MatB>::size1() const [with MatA=remora::dense_triangular_proxy<const float, remora::row_major, remora::lower, remora::hip_tag>, MatB=remora::matrix<float, remora::column_major, remora::hip_tag>]" 
/home/user/Remora/include/remora/cpu/../assignment.hpp(258): here
            instantiation of "MatA &remora::assign(remora::matrix_expression<MatA, Device> &, const remora::matrix_expression<MatB, Device> &) [with MatA=remora::dense_matrix_adaptor<float, remora::row_major, remora::continuous_dense_tag, remora::hip_tag>, MatB=remora::matrix_matrix_prod<remora::dense_triangular_proxy<const float, remora::row_major, remora::lower, remora::hip_tag>, remora::matrix<float, remora::column_major, remora::hip_tag>>, Device=remora::hip_tag]" 
/home/user/Remora/include/remora/cpu/../assignment.hpp(646): here
            instantiation of "remora::noalias_proxy<C>::closure_type &remora::noalias_proxy<C>::operator=(const E &) [with C=remora::matrix<float, remora::row_major, remora::hip_tag>, E=remora::matrix_matrix_prod<remora::dense_triangular_proxy<const float, remora::row_major, remora::lower, remora::hip_tag>, remora::matrix<float, remora::column_major, remora::hip_tag>>]" 
/home/user/Remora/Test/hip_triangular_prod.cpp(325): here
            instantiation of "void Remora_hip_triangular_prod::triangular_prod_matrix_matrix_test(Orientation) [with Orientation=remora::row_major]" 
/home/user/Remora/Test/hip_triangular_prod.cpp(527): here

Ответы [ 3 ]

0 голосов
/ 04 сентября 2018

Решение, которое я придумал в то время с гораздо меньшим дублированием кода, состоит в том, чтобы заменить вызов уровнем функтора: 

template<class Impl, class Device>
struct WrapperImpl;
template<class Impl>
struct WrapperImpl<Impl, CPU>{
    typename Impl::Functor f;
    __host__ operator()(){ f();}
};
//identical to CPU up to __device__
template<class Impl>
struct WrapperImpl<Impl, GPU>{
    typename Impl::Functor f;
    __device__ operator()(){ f();}
};

template<class Impl>
struct Wrapper{
    typedef WrapperImpl<Impl, typename Impl::Device> Functor;
    Impl impl;
    // lots and lots of decorator code that i now do not need to duplicate
    Functor call_functor()const{
        return Functor{impl.call_functor();};
    }
};

//repeat for around 20 classes

Wrapper<ImplCPU> wrapCPU;
wrapCPU.call_functor()();
0 голосов
/ 06 сентября 2018

Эта проблема на самом деле является довольно прискорбным недостатком языковых расширений CUDA.

Стандартный подход для обработки этих предупреждений (в Thrust и аналогичных шаблонных библиотеках CUDA) заключается в отключении предупреждения для функции / метода, вызывающего его, с помощью #pragma hd_warning_disable или в более новой CUDA (9.0 или новее) #pragma nv_exec_check_disable .

Так что в вашем случае это будет: 

template<class Impl>
struct Wrapper{
    Impl impl;
    // lots and lots of decorator code

      #pragma nv_exec_check_disable
    __host__ __device__ void call(){ impl.call();};
};

Подобный вопрос уже задавался

0 голосов
/ 04 сентября 2018

Извините, но вы злоупотребляете языком и вводите в заблуждение читателей. Неверно, что у ваших классов-обёрток есть метод __host__ __device__; Вы хотите сказать, что он имеет или метод __host__ или метод __device__. Вы должны рассматривать предупреждение как ошибку.

Таким образом, вы не можете просто использовать пример шаблона для ImplCPU и ImplGPU; но - ты мог бы сделать что-то подобное? 

template<typename Impl> struct Wrapper;

template<> struct Wrapper<ImplGPU> {
    ImplGPU impl;
    __device__ void call(){ impl.call();};
}

template<> struct Wrapper<ImplCPU> {
    ImplGPU impl;
    __host__ void call(){ impl.call();};
}

или если вы хотите быть более педантичным, это может быть: 

enum implementation_device { CPU, GPU };

template<implementation_device ImplementationDevice> Wrapper;
template<> Wrapper<CPU> {
    __host__ void call();
}
template<> Wrapper<GPU> {
    __device__ void call();
}

Сказав это - вы ожидали использовать один класс Wrapper, и здесь я говорю вам, что вы не можете этого сделать. Я подозреваю, что ваш вопрос представляет собой проблему X-Y , и вы действительно должны рассмотреть весь подход к использованию этой оболочки. Возможно, вам нужен код, который использует его по-разному для процессора или графического процессора. Возможно, вам нужно где-нибудь удалить тип. Но это не сработает.

...