CUDA中不允許如何修復調用__host__函數的錯誤（「std :: max <double>」）？

Question

#include <algorithm> 
#include <vector> 
template <typename Dtype> 
    __global__ void R_D_CUT(const int n, Dtype* r, Dtype* d 
     , Dtype cur_r_max, Dtype cur_r_min, Dtype cur_d_max, Dtype cur_d_min) { 
     CUDA_KERNEL_LOOP(index, n) { 
      r[index] = __min(cur_r_max, __max(r[index], cur_r_min)); 
      d[index] = __min(cur_d_max, __max(d[index], cur_d_min)); 
     } 
    } 

在上面的代碼中，它可以在Window中很好地工作。但是，由於__min和__max函數，它在Ubuntu中不起作用。要修復它通過更換__min到std::min<Dtype>和max到std::max<Dtype>：

template <typename Dtype> 
    __global__ void R_D_CUT(const int n, Dtype* r, Dtype* d 
     , Dtype cur_r_max, Dtype cur_r_min, Dtype cur_d_max, Dtype cur_d_min) { 
     CUDA_KERNEL_LOOP(index, n) { 

      r[index] = std::min<Dtype>(cur_r_max, std::max<Dtype>(r[index], cur_r_min)); 
      d[index] = std::min<Dtype>(cur_d_max, std::max<Dtype>(d[index], cur_d_min)); 
     } 
    } 

然而，當我重新編譯，我得到了錯誤

_layer.cu(7): error: calling a __host__ function("std::min<float> ") from a __global__ function("caffe::R_D_CUT<float> ") is not allowed 

_layer.cu(7): error: calling a __host__ function("std::max<float> ") from a __global__ function("caffe::R_D_CUT<float> ") is not allowed 

_layer_layer.cu(8): error: calling a __host__ function("std::min<float> ") from a __global__ function("caffe::R_D_CUT<float> ") is not allowed 

_layer_layer.cu(8): error: calling a __host__ function("std::max<float> ") from a __global__ function("caffe::R_D_CUT<float> ") is not allowed 

_layer_layer.cu(7): error: calling a __host__ function("std::min<double> ") from a __global__ function("caffe::R_D_CUT<double> ") is not allowed 

_layer_layer.cu(7): error: calling a __host__ function("std::max<double> ") from a __global__ function("caffe::R_D_CUT<double> ") is not allowed 

_layer_layer.cu(8): error: calling a __host__ function("std::min<double> ") from a __global__ function("caffe::R_D_CUT<double> ") is not allowed 

_layer_layer.cu(8): error: calling a __host__ function("std::max<double> ") from a __global__ function("caffe::R_D_CUT<double> ") is not allowed 

你能幫我解決嗎？謝謝

Answer 1

一般來說，在CUDA設備代碼（__global__或__device__函數）中，與std::關聯的功能不可用。

相反，對於許多數學功能，NVIDIA提供了CUDA math library。

對於這種情況，正如@njuffa指出的那樣，CUDA提供了min和max的模板/過載版本。所以你應該只能在設備代碼中使用min()或max()，假設類型使用對應於可用的模板/重載類型之一。此外，你應該：

#include <math.h> 

這裏是顯示的min()使用一個簡單的工作實例都float和double類型：

$ cat t381.cu 
#include <math.h> 
#include <stdio.h> 

template <typename T> 
__global__ void mymin(T d1, T d2){ 

    printf("min is :%f\n", min(d1,d2)); 
} 


int main(){ 

    mymin<<<1,1>>>(1.0, 2.0); 
    mymin<<<1,1>>>(3.0f, 4.0f); 
    cudaDeviceSynchronize(); 
} 
$ nvcc -arch=sm_52 -o t381 t381.cu 
$ ./t381 
min is :1.000000 
min is :3.000000 
$ 

注意，可用過載選項甚至include some integer types

Answer 2

添加到@ RobertCrovella的answer：如果您想要的行爲更像std::max，您可以在CUDA的數學庫上使用此模板化包裝：

#define __df__ __device__ __forceinline__ 
template <typename T> __df__ T maximum(T x, T y); 
template <> __df__ int     maximum<int    >(int x, int y)        { return max(x,y); } 
template <> __df__ unsigned int  maximum<unsigned   >(unsigned int x, unsigned int y)    { return umax(x,y); } 
template <> __df__ long    maximum<long    >(long x, long y)        { return llmax(x,y); } 
template <> __df__ unsigned long  maximum<unsigned long  >(unsigned long x, unsigned long y)   { return ullmax(x,y); } 
template <> __df__ long long   maximum<long long   >(long long x, long long y)     { return llmax(x,y); } 
template <> __df__ unsigned long long maximum<unsigned long long>(unsigned long long x, unsigned long long y) { return ullmax(x,y); } 
template <> __df__ float    maximum<float    >(float x, float y)       { return fmaxf(x,y); } 
template <> __df__ double    maximum<double   >(double x, double y)       { return fmax(x,y); } 
#undef __df__ 

_{（見here更完整的一套這些包裝的。）}