0

我正在 CUDA 中做一些原子实验。我的大问号是，当在同一个块中运行的两个线程以原子方式访问同一个地址时，它们的行为如何。我用 atomicAdd 进行了一些测试，它以原子方式工作，但是当我使用 atomicCAS 尝试下面的代码时，结果不是我所期望的。有人有解释吗？

#include <cuda_runtime.h>
#include <iostream>
#include <cuComplex.h>
using namespace std;
__global__ void kernel(int * pointer)
{
    *pointer=0;
    *(pointer+threadIdx.x+1)=0;
    __syncthreads();
    *(pointer+threadIdx.x+1)=atomicCAS(pointer,0,100);
}
int main(int argc,char ** argv)
{
    int numThreads=40;
    dim3 threadsPerBlock;
    dim3 blocks;
    int o[numThreads+1];
    int * pointer;
    cudaMalloc(&pointer,sizeof(int)*(numThreads+1));
    cudaMemset(pointer,0,sizeof(int)*(numThreads+1));
    threadsPerBlock.x=numThreads;
    threadsPerBlock.y=1;
    threadsPerBlock.z=1;
    blocks.x=1;
    blocks.y=1;
    blocks.z=1;
    kernel <<<threadsPerBlock,blocks>>> (pointer);
    cudaMemcpy(o,pointer,sizeof(int)*(numThreads+1),cudaMemcpyDeviceToHost);



    for (int i=0;i<numThreads+1;i++)
            cout << o[i] << " ";

    cout << endl;

}

在上面的代码 atomicCAS 在同一块内运行访问相同的地址以进行比较和交换......我的期望是只有一个 atomicCAS 会找到与 0 比较的值，而所有其他人会发现它 100，但奇怪的是输出我的程序是：

 100 100 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0

即所有线程都找到要比较的值设置为0。

4

2 回答 2

5

您已经颠倒了执行配置参数的顺序。是<<<gridDim, blockDim>>>，反之亦然。所以你要启动 40 个块，每个块有 1 个线程，而不是相反。

这就是为什么你会得到你看到的结果——因为每个块中只有一个线程在运行，numThreads-1所以数组中的最后一个值总是为零。

如果我交换订单，我会得到这个输出：

100 100 100 100 100 100 100 100 100 100 100 100 100 100 100 100 100 100 100 100 100 100 100 100 100 100 100 100 100 100 100 100 100 0 100 100 100 100 100 100 100

您可以看到除一个线程外的所有线程都按预期写入100，并且一个线程已写入。0

于 2012-11-22T23:21:35.117 回答
2

你有你的线程块并在你的内核调用中向后阻塞变量。

而不是这个：

 kernel <<<threadsPerBlock,blocks>>> (pointer);

做这个：

 kernel <<<blocks, threadsPerBlock>>> (pointer);

然后你会得到正确的输出。

于 2012-11-22T23:34:20.567 回答