nvcc предупреждает о том, что переменная устройства является переменной хоста - почему? - PullRequest
Новая
       53

nvcc предупреждает о том, что переменная устройства является переменной хоста - почему?

1 голос
/ 28 марта 2019

Я читал в Руководстве по программированию CUDA о функциях шаблонов и работает ли что-то подобное?

#include <cstdio>

/* host struct */
template <typename T>
struct Test {
    T  *val;
    int size;
};

/* struct device */
template <typename T>
__device__ Test<T> *d_test;

/* test function */
template <typename T>
T __device__ testfunc() {
    return *d_test<T>->val;
}

/* test kernel */
__global__ void kernel() {
    printf("funcout = %g \n", testfunc<float>());
}

Я получаю правильный результат, но предупреждение:

"предупреждение: переменная хоста" d_test [with T = T] "не может быть непосредственно прочитана в функции устройства"?

Имеет структуру в тестовой функции для создания экземпляра с *d_test<float>->val?

KR, Iggi

Ответы [ 2 ]

2 голосов
/ 28 марта 2019

К сожалению, компилятор CUDA, похоже, обычно имеет некоторые проблемы с переменными шаблонами. Если вы посмотрите на сборку , вы увидите, что все работает просто отлично. Компилятор явно создает экземпляр шаблона переменной и выделяет соответствующий объект устройства. 

.global .align 8 .u64 _Z6d_testIfE;

Сгенерированный код использует этот объект так же, как он должен 

ld.global.u64   %rd3, [_Z6d_testIfE];

Я бы посчитал это предупреждение ошибкой компилятора. Обратите внимание, что я не могу воспроизвести проблему с CUDA 10 здесь, поэтому эта проблема, скорее всего, уже исправлена. Подумайте об обновлении вашего компилятора & hellip;

1 голос
/ 28 марта 2019

@ MichaelKenzel верен.

Это почти наверняка ошибка nvcc - которую я сейчас подал (вам может понадобиться учетная запись для доступа к ней.

Такжепримечание: я смог воспроизвести проблему с меньшим количеством кода: 

template <typename T>
struct foo { int  val; };

template <typename T>
__device__ foo<T> *x;

template <typename T>
int __device__ f() { return x<T>->val; }

__global__ void kernel() { int y = f<float>(); }

, а также посмотреть на результат на GodBolt .

...