Mẫu __host__ __device__ gọi các chức năng được xác định bởi máy chủ

Question

Khi triển khai mã CUDA, tôi thường cần một số chức năng tiện ích, được gọi từ thiết bị và cũng từ mã máy chủ. Vì vậy, tôi khai báo các hàm này là __host__ __device__. Điều này là ổn và có thể xảy ra do thiết bị/máy chủ lưu trữ có thể được xử lý bởi #ifdef CUDA_ARCH.

Sự cố xảy ra khi chức năng tiện ích được tạo mẫu tức là. bởi một số loại functor. Nếu mẫu Ví dụ gọi một hàm __host__ tôi nhận được cảnh báo này:

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

Chỉ giải pháp tôi biết là để xác định chức năng hai lần - một lần cho thiết bị và một lần cho mã máy chủ với tên khác (Tôi không thể quá tải trên __host__ __device__). Nhưng điều này có nghĩa rằng có sự sao chép mã và tất cả các chức năng khác của __host__ __device__ sẽ gọi nó, cũng phải được xác định hai lần (thậm chí nhiều hơn việc sao chép mã).

ví dụ đơn giản:

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

Cảnh báo là do foo(HostObject()); gọi bên trong main() chức năng.

foo_host<> và foo_device<> có thể thay thế cho sự cố foo<>.

Có giải pháp nào tốt hơn không? Tôi có thể ngăn chặn hiện tượng foo() ở phía thiết bị không?

Answer 1

Bạn không thể ngăn việc khởi tạo một nửa số phiên bản mẫu chức năng __host__ __device__. Nếu bạn khởi tạo hàm bằng cách gọi nó trên máy chủ (thiết bị), trình biên dịch cũng sẽ khởi tạo thiết bị (máy chủ) một nửa.

Điều tốt nhất bạn có thể làm cho trường hợp sử dụng của mình là CUDA 7.0 là chặn cảnh báo bằng cách sử dụng #pragma hd_warning_disable như trong ví dụ sau và đảm bảo rằng chức năng không được gọi sai.

#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; 
}