使用 cos 时，Cuda 从 __device__ 函数返回错误值

Question

我正在尝试在 nvidia GPU 上运行此代码，但它返回奇怪的值。它由 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，它返回正确 结果。
• 当我在exmodul.cu 中更改1.1 时，除了值1.0 之外，它也会失败。
• 行为不依赖于同一模块中的常量2.0。
• 当我使用另一个函数而不是 cos（如 sin 或 exp）时，它可以正常工作。
• 使用double q = cos(1.1); 无效。
• 当我将函数 extFunc() 复制到模块 main.cu 时，它可以正常工作。
• 如果我在main.cu 中取消注释*gd=1.0;，它会返回正确的1.0。

在 Nvidia GT750M 和 GeForce GTX TITAN Black 上测试。 （在 GT750M 上返回不同的值 gd=6.1232329394368592e-17 但仍然错误）。操作系统：Debian Jessie。

$ 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 来解决这个问题。我怀疑编译过程中存在一个已经修复的潜在错误。