1

我在程序中有一个非常奇怪的错误。我花了很多时间在上面，但我还没有找到解决方案。我编写了简单的程序来重现我的问题。也许有人帮助我。我尝试了 cuda-memcheck &使用 CUDA 运行时 API 检查错误的规范方法是什么？但我没有收到任何错误。

细节：

nvcc 版本 - V6.0.1

gcc 版本 - 4.8.1

完整代码：

#include <stdio.h>

__constant__ unsigned long long int bigNumber = 83934243334343;
__device__ bool isFound = false;
__global__ void kernel(int *dev_number) {

    unsigned long long int id = threadIdx.x + (blockIdx.x * blockDim.x);
    while (id < bigNumber && isFound==false) {

        if(id == 10) {
            *dev_number = 4;
            isFound=true;
        }
        id++;
    }
}

int main(int argc, char *argv[]) {
    int number = 0;
    int *dev_number;

    printf("Number: %d\n", number);

    return 0;
}

编译并运行：

nvcc myprogram.cu
./myprogram

当我运行这个程序时，我没有得到任何返回值。但是当变量 - bigNumber 具有较小的值或者我不使用 cudaMalloc 和 cudaMemcpy 时它可以工作（这意味着调用 return 0）。什么连接必须为另一个具有常量 bigNumber 的变量分配内存？有什么问题？

4

1 回答 1

1

现在您已将代码修改为更合理的内容，我通过以下修改立即获得结果：

__device__ volatile bool isFound = false;

volatile限定符强制编译器忽略任何会阻止每个线程读取变量的全局副本的优化。

从文档

编译器可以自由优化对全局或共享内存的读取和写入（例如，通过将全局读取缓存到寄存器或 L1 缓存中），只要它尊重内存围栏函数（Memory Fence Functions）的内存排序语义和内存可见性语义同步函数（Synchronization Functions）。

可以使用 volatile 关键字禁用这些优化：如果将位于全局或共享内存中的变量声明为 volatile，编译器假定它的值可以随时被另一个线程更改或使用，因此对该变量的任何引用都会编译为实际的内存读取或写入指令。

如果你没有使用volatile限定符，那么只有一个线程采用提前退出条件（isFound），所有其他线程必须循环很长时间，直到它们的id值超过bigNumber

于 2015-04-13T14:12:07.600 回答