2015-05-04 65 views
9

在執行CUDA代碼的過程中,我經常需要一些實用函數,它們將從設備和主機代碼中調用。所以我聲明這些函數爲__host__ __device__。這是可以的,並且可能的設備/主機不兼容性可以通過#ifdef CUDA_ARCH來處理。模板__host__ __device__調用主機定義的函數

當效用函數模板化時出現問題,由一些函子類型。如果模板實例調用一個__host__功能我得到這樣的警告:

calling a __host__ function from a __host__ __device__ function is not allowed 
     detected during instantiation of "int foo(const T &) [with T=HostObject]" 

我知道唯一的辦法就是兩次定義函數 - 一次設備和一次用不同的名稱主機代碼(我不能__host__ __device__超載)。但這意味着代碼重複,並且所有其他將調用它的函數也必須定義兩次(甚至更多的代碼重複)。

簡化的例子:

#include <cuda.h> 
#include <iostream> 

struct HostObject { 
    __host__ 
    int value() const { return 42; } 
}; 

struct DeviceObject { 
    __device__ 
    int value() const { return 3; } 
}; 

template <typename T> 
__host__ __device__ 
int foo(const T &obj) { 
    return obj.value(); 
} 

/* 
template <typename T> 
__host__ 
int foo_host(const T &obj) { 
    return obj.value(); 
} 

template <typename T> 
__device__ 
int foo_device(const T &obj) { 
    return obj.value(); 
} 
*/ 

__global__ void kernel(int *data) { 
    data[threadIdx.x] = foo(DeviceObject()); 
} 

int main() { 
    foo(HostObject()); 

    int *data; 
    cudaMalloc((void**)&data, sizeof(int) * 64); 
    kernel<<<1, 64>>>(data); 
    cudaThreadSynchronize(); 
    cudaFree(data); 
} 

警告由main()函數內部的foo(HostObject());呼叫造成的。

foo_host<>foo_device<>可能替代有問題的foo<>

有沒有更好的解決方案?我可以在設備端防止foo()的實例嗎?

+0

沒有構造函數在'foo()'中被調用。問題正是警告所說的。我問我是否可以在不定義泛型函數兩次的情況下修正它。 – Johny

+0

該警告是由主函數中的foo(HostObject())引起的。構造函數沒有問題,因爲除非我自己聲明一個,否則會自動生成構造函數(由主機和設備編譯器)。 – Johny

+0

對不起,現在我明白了你的觀點 - 不容易看出沒有編譯器的情況下顯示錯誤的位置。因此,在我相信的問題中提及它會很有用。 –

回答

5

您無法阻止實例化__host__ __device__函數模板實例的一半。如果通過在主機(設備)上調用該函數來實例化該函數,編譯器也將實例化設備(主機)的一半。

對於您的CUDA 7.0用例,您可以做的最好的做法是使用#pragma hd_warning_disable來抑制警告,並確保該函數未被錯誤地調用。

#include <iostream> 
#include <cstdio> 

#pragma hd_warning_disable 
template<class Function> 
__host__ __device__ 
void invoke(Function f) 
{ 
    f(); 
} 

struct host_only 
{ 
    __host__ 
    void operator()() 
    { 
    std::cout << "host_only()" << std::endl; 
    } 
}; 

struct device_only 
{ 
    __device__ 
    void operator()() 
    { 
    printf("device_only(): thread %d\n", threadIdx.x); 
    } 
}; 

__global__ 
void kernel() 
{ 
    // use from device with device functor 
    invoke(device_only()); 

    // XXX error 
    // invoke(host_only()); 
} 

int main() 
{ 
    // use from host with host functor 
    invoke(host_only()); 

    kernel<<<1,1>>>(); 
    cudaDeviceSynchronize(); 

    // XXX error 
    // invoke(device_only()); 

    return 0; 
}