टेम्पलेट __host__ __device__ कॉलिंग होस्ट परिभाषित फ़ंक्शंस

Question

सीयूडीए कोड के कार्यान्वयन के दौरान मुझे अक्सर कुछ उपयोगिता कार्यों की आवश्यकता होती है, जिन्हें डिवाइस से और होस्ट कोड से भी कहा जाएगा। तो मैं इन कार्यों को __host__ __device__ के रूप में घोषित करता हूं। यह ठीक है और संभव डिवाइस/मेजबान अक्षमता #ifdef CUDA_ARCH द्वारा संभाला जा सकता है।

समस्याएं तब आती हैं जब उपयोगिता फ़ंक्शन टेम्पलेट होता है यानी। कुछ मज़ेदार प्रकार से। - एक बार डिवाइस के लिए और एक बार भिन्न नाम के साथ मेजबान कोड के लिए (मैं __host__ __device__ पर ओवरलोड नहीं कर सकते

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

केवल समाधान मैं जानता हूँ कि समारोह में दो बार परिभाषित करने के लिए है: टेम्पलेट उदाहरण एक __host__ फ़ंक्शन को कॉल यदि मैं यह चेतावनी मिलती है)। लेकिन इसका मतलब है कि कोड डुप्लिकेशन और अन्य सभी __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() के इंस्टॉलेशन को रोक सकता हूं?

Answer 1

आप __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; 
}