Статическая переменная __device__ и ядра в отдельном файле

Question

Я хочу статически объявить глобальную переменную с квалификатором __device__. В то же время я хочу хранить функции, предназначенные для GPU, в отдельном файле.

Однако, если я это сделаю, значение переменной не передается в GPU - нет ошибок во время компиляции или выполнения, но функции memcpy ничего не делают.

Когда я перемещаю функцию ядрав файл с кодом хоста все работает.

Я уверен, что в этом случае должна быть возможность разбить функции хоста и устройства на отдельные файлы, но как это сделать? Я видел только примеры, когда ядра и код хоста находятся в одном файле.

Я также был бы очень благодарен, если бы кто-то объяснил, почему он так себя ведет.

Пример кодаперечислены ниже.

Заранее спасибо.

---

Рабочий каталог:

$ ls
functionsGPU.cu  functionsGPU.cuh  staticGlobalMemory.cu

staticGlobalMemory.cu:

#include "functionsGPU.cuh"

#if VARIANT == 2
__global__ void checkGlobalVariable(){
    printf("Old value (dev): %f\n", devData);
    devData += 2.0f;
    printf("New value (dev): %f\n", devData);
}
#endif

int main(int argc, char **argv){
    int dev = 0;
    float val = 3.2;

    cudaSetDevice(dev);

    printf("---------\nVARIANT %i\n---------\n", VARIANT);

    printf("Old value (host): %f\n", val);
    cudaMemcpyToSymbol(devData, &val, sizeof(float));
    checkGlobalVariable <<<1, 1>>> ();
    cudaMemcpyFromSymbol(&val, devData, sizeof(float));
    printf("New value (host): %f\n", val);

    cudaDeviceReset();

    return 0;
}

functionsGPU. cuh:

#ifndef FUNCTIONSGPU_CUH
#define FUNCTIONSGPU_CUH

#include <cuda_runtime.h>
#include <stdio.h>

#define VARIANT 1

__device__ float devData;

#if VARIANT == 1
__global__ void checkGlobalVariable();
#endif

#endif

functionsGPU.cu:

#include "functionsGPU.cuh"

#if VARIANT == 1
__global__ void checkGlobalVariable(){
    printf("Old value (dev): %f\n", devData);
    devData += 2.0f;
    printf("New value (dev): %f\n", devData);
}
#endif

Это скомпилировано как

$ nvcc -arch=sm_61 staticGlobalMemory.cu functionsGPU.cu -o staticGlobalMemory

Вывод, если код ядра и хоста находятся в отдельных файлах (неверно):

---------
VARIANT 1
---------
Old value (host): 3.200000
Old value (dev): 0.000000
New value (dev): 2.000000
New value (host): 3.200000

Выводится, если код ядра и хоста находится в одном файле (правильно):

---------
VARIANT 2
---------
Old value (host): 3.200000
Old value (dev): 3.200000
New value (dev): 5.200000
New value (host): 5.200000

Answer 1

Ваша структура кода, где код устройства в одном модуле компиляции ссылается на код устройства или объекты устройства в другом модуле компиляции, потребует компиляции и связывания кода перемещаемого устройства CUDA.

1. Для включения -rdc=trueэто, к вашей nvcc командной строке компиляции
2. Добавьте extern перед определением devData, в functionsGPU.cuh
3. Добавьте __device__ float devData; к staticGlobalMemory.cu

Это должно решить проблему. Шаг 1 обеспечивает необходимую межмодульную связь, а шаги 2 и 3 устранят проблему дублирования определений, которая у вас возникнет, поскольку вы включаете одну и ту же переменную через файл заголовка в отдельные блоки компиляции.