使用 cos 时，Cuda returns 来自 __device__ 函数的错误值

Question

我正尝试在 nvidia GPU 上 运行 这段代码，它 return 的值很奇怪。它由两个模块 main.cu 和 exmodul.cu 组成（ 如下所列 ）。对于建筑，我正在使用：

nvcc -dc -arch sm_35 main.cu
nvcc -dc -arch sm_35 exmodul.cu
nvcc     -arch sm_35 -lcudart -o main main.o exmodul.o

如果我运行那我最后一行就搞怪了！！！ gd 必须为 1.

result=0
result=0
result=0
result=0
gd=-0.5

• 当我将 exmodul.cu 中的 1.0 更改为大于 1.000000953 或以下 0.999999999999999945, return 正确 结果。
• 当我在 exmodul.cu 中更改 1.1 时，它也失败了，除了值 1.0.
• 行为不依赖于同一模块中的常量 2.0。
• 当我使用另一个函数而不是 cos 时，如 sin 或 exp 它工作正常。
• 使用 double q = cos(1.1); 无效。
• 当我将函数 extFunc() 复制到模块 main.cu 时，它工作正常。
• 如果我在 main.cu 中取消注释 *gd=1.0; 它 return 是正确的 1.0。

在 Nvidia GT750M 和 GeForce GTX TITAN Black 上测试。 （在 GT750M 上 return 的不同值 gd=6.1232329394368592e-17 但仍然错误）。 OS：Debian 杰西。

$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2013 NVIDIA Corporation
Built on Thu_Mar_13_11:58:58_PDT_2014
Cuda compilation tools, release 6.0, V6.0.1

你知道哪里出了问题吗？

谢谢，卢卡斯

---

main.cu

#include <stdio.h>   // printf
#include "exmodul.h" // extFunc

__global__ void mykernel(double*gd);
void deviceCheck();

int main(int argc, char *argv[])
{
    double gd, *d_gd;
    cudaMalloc(&d_gd, sizeof(double)); deviceCheck();
    mykernel<<<1,1>>>(d_gd);           deviceCheck();
    cudaMemcpy(&gd, d_gd, sizeof(double), cudaMemcpyDeviceToHost);
                                       deviceCheck();
    cudaFree(d_gd);                    deviceCheck();
    fprintf(stderr,"gd=%.17g\n",gd);
    return 0;
}

void deviceCheck()
{
    cudaError_t result = cudaSuccess;
    cudaDeviceSynchronize();
    result = cudaGetLastError();
    fprintf(stderr,"result=%d\n",result); fflush(stderr);
}

__global__ void mykernel(double *gd)
{
    *gd = extFunc();
    //*gd=1.0;
    __syncthreads();
    return;
}

exmodul.cu

#include "exmodul.h"

__device__ double extFunc()
{
    double q = 1.1;
    q = cos(q);
    if(q<2.0) { q = 1.0; }
    return q;
}

exmodul.h

__device__ double extFunc();

Answer 1

我能够在受支持的配置（CUDA 6.5、CentOS 6.2、K40）上重现有问题的行为。

当我从 CUDA 6.5 切换到 CUDA 7 RC 时，问题消失了。

该问题似乎也无法在较旧的配置（CUDA 5.5、CentOS 6.2、M2070）上重现

我建议切换到 CUDA 7 RC 来解决这个问题。我怀疑编译过程中的潜在错误已经修复。