Использование функционально-шаблонного кода через границу g ++ — nvcc (включая ядра)
Предположим, я скомпилировал следующее с помощью nvcc компилятора NVIDIA CUDA:
template<typename T, typename Operator>
__global__ void fooKernel(T t1, T t2) {
Operator op;
doSomethingWith(t1, t2);
}
template<typename T>
__device__ __host__ void T bar(T t1, T t2) {
return t1 + t2;
}
template<typename T, typename Operator>
void foo(T t1, T t2) {
fooKernel<<<2, 2>>>(t1, t2);
}
// explicit instantiation
template decltype(foo<int, bar<int>>) foo<int, bar<int>);
Теперь я хочу, чтобы мой gcc, не nvcc код вызывал foo():
...
template<typename T, typename Operator> void foo(T t1, T t2);foo<int, bar<int>> (123, 456);
...
У меня есть соответствующий (?) Экземпляр в файле .o / .a / .so, который я компилирую с CUDA.
Могу ли я сделать эту работу как-нибудь, несмотря на смешение __device__ функция с tr __nondevice__ один?
Решение
Проблема здесь в том, что шаблонный код обычно создается в месте использования, что не работает, потому что foo содержит вызов ядра, который не может быть проанализирован g ++. Ваш подход к явной реализации шаблона и его объявлению для компилятора хоста является правильным. Вот как это сделать. Я немного исправил ваш код и разделил его на 3 файла:
- gpu.cu
- gpu.cuh
- cpu.cpp
gpu.cuh
Этот файл содержит шаблонный код для использования gpu.cu, Я добавил какую-то цель к твоему foo() функция, чтобы убедиться, что это работает.
#pragma once
#include <cuda_runtime.h>
template <typename T>
struct bar {
__device__ __host__ T operator()(T t1, T t2)
{
return t1 + t2;
}
};
template <template <typename> class Operator, typename T>
__global__ void fooKernel(T t1, T t2, T* t3)
{
Operator<T> op;
*t3 = op(t1, t2);
}
template <template <typename> class Operator, typename T>
T foo(T t1, T t2)
{
T* t3_d;
T t3_h;
cudaMalloc(&t3_d, sizeof(*t3_d));
fooKernel<Operator><<<1, 1>>>(t1, t2, t3_d);
cudaMemcpy(&t3_h, t3_d, sizeof(*t3_d), cudaMemcpyDeviceToHost);
cudaFree(t3_d);
return t3_h;
}
gpu.cu
Этот файл только создает foo() Функция, чтобы убедиться, что она будет доступна для ссылки:
#include "gpu.cuh"
template int foo<bar>(int, int);
cpu.cpp
В этом простом исходном файле C ++ нам нужно убедиться, что мы не получаем экземпляры шаблона, так как это приведет к ошибке компиляции. Вместо этого мы только вперед объявляем структуру bar и функция foo, Код выглядит так:
#include <cstdio>
template <template <typename> class Operator, typename T>
T foo(T t1, T t2);
template <typename T>
struct bar;
int main()
{
printf("%d \n", foo<bar>(3, 4));
}
Makefile
Это соберет весь код в исполняемый файл:
.PHONY: clean all
all: main
clean:
rm -f *.o main
main: gpu.o cpu.o
g++ -L/usr/local/cuda/lib64 $^ -lcudart -o $@
gpu.o: gpu.cu
nvcc -c -arch=sm_20 $< -o $@
cpu.o: cpu.cpp
g++ -c $< -o $@
Код устройства компилируется nvccкод хоста g++ и все это связано g++, После запуска вы видите прекрасный результат:
7
Здесь следует помнить, что запуск ядра и определения ядра должны быть в .cu файлы, которые компилируются nvcc, Для дальнейшего использования я также оставлю эту ссылку здесь, на разделение ссылок и компиляции с CUDA.
Другие решения