2

Suppose I compile the following with NVIDIA CUDA's nvcc compiler:

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>);

Now, I want my gcc, non-nvcc code to call foo():

...

template<typename T, typename Operator> void foo(T t1, T t2);


foo<int, bar<int>> (123, 456);
...

I have the appropriate (?) instantiation in the .o/.a/.so file I compile with CUDA.

Can I make that happen?

4

1 回答 1

2

这里的问题是模板化代码通常在使用位置进行实例化，这不起作用，因为foo包含无法被 g++ 解析的内核调用。您显式实例化模板并为主机编译器转发声明它的方法是正确的。这是如何做到这一点的。我稍微修正了您的代码并将其拆分为 3 个文件：

  1. 图形处理器
  2. 显卡
  3. cpu.cpp

显卡

此文件包含供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;
}

图形处理器

此文件仅实例化该foo()函数以确保它可用于链接：

#include "gpu.cuh"

template int foo<bar>(int, int);

cpu.cpp

在这个普通的 C++ 源文件中，我们需要确保我们没有获得模板实例化，因为这会产生编译错误。相反，我们只转发声明 structbar和 function 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));
}

生成文件

这会将代码全部放在一个可执行文件中：

.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 的链接和编译分离。

于 2015-02-09T19:09:17.900 回答