如何修复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

一般而言,与std::CUDA设备代码(__global__或__device__函数)中不可用的功能相关.

相反,对于许多数学函数,NVIDIA提供了一个CUDA数学库.

对于这种情况,正如@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
$

请注意,可用的重载选项甚至包括一些整数类型