2

NVIDIA CUDA の nvcc コンパイラで以下をコンパイルするとします。

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

CUDA でコンパイルした .o/.a/.so ファイルに適切な (?) インスタンス化があります。

私はそれを実現できますか？

4

1 に答える 1

2

ここでの問題は、テンプレート化されたコードは通常、使用場所でインスタンス化されますがfoo、g++ で解析できないカーネル呼び出しが含まれているため機能しません。テンプレートを明示的にインスタンス化し、それをホスト コンパイラに対して前方宣言するというアプローチは正しいものです。これを行う方法は次のとおりです。コードを少し修正し、3 つのファイルに分割しました。

  1. gpu.cu
  2. gpu.cuh
  3. 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));
}

メイクファイル

これにより、コードがすべて実行可能ファイルにまとめられます。

.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 に答える