c++ - 如何使用threadfence/CUDA5.5

Question

所以，我想知道如何使用threadfence，我想阅读关于threadfence的示例代码。

请告诉我有关 threadfence() 的示例代码

我相信当我想访问内存是CUDA5.5中设备端的内核函数时，我想执行独占控制。
我了解到，如果我使用所谓的“() __threadfence”，则可以进行独占控制，但我不知道如何使用。

· 我能做些什么？
(现在，“__threadfence()”发生了未定义标识的错误)
· 我在源代码上的什么地方写代码。 因为我想我想在下面的代码中独占控制要访问 Log_d 的地方。

其他线程我想阻止在“锁定开始”~“锁定停止”之间访问 Log_d 的内存，例如以下代码。

我在 CPU 端代码中定义了 block 和线程。 block :1,1,1 和线程:256,1,1

__global__ void matrix_vector_multi_gpu_1_256(float *A_d, float *B_d, float *C_d, float *Log_d){
    int i;

    A_d[threadIdx.x]=0.0F;

    for(i=0;i<N;i++){
        A_d[threadIdx.x]=A_d[threadIdx.x]+B_d[threadIdx.x*N+i]*C_d[i];
    }
    //lock Start about Log_d
    //__threadfence();
    for(int j=0;j<N;j++){
        if(Log_d[j]==0){
            Log_d[j]=threadIdx.x + 1;
            break;
        }
    }
    //Stop the lock


}

Answer 1

如果您能提供更多关于matrix_vector_multi_gpu_1_256 内核函数应该做什么的信息，那就太好了。

通过 __threadfence()，设备会一直等待，直到调用线程进行的所有全局和共享访问对以下对象可见:

1. 线程 block 中的所有线程用于共享内存访问；
2. 设备中用于全局内存访问的所有线程。

CUDA SDK 的 threadFenceReduction 中提供了使用 __threadfence() 的示例。

在示例中，缩减是在单个内核调用中对任意大小的数组执行的。线程 block 执行部分缩减，内核通过全局计数器的原子增量跟踪完成了多少 block 。如果票的值等于线程 block 的数量，那么持有票的 block 就知道它是最后一个要完成的 block 。最后一个 block 负责对所有其他 block 的结果求和。

为了使这种方法正常工作，必须确保在区 block 获取票证之前，其所有内存事务都已完成。这是通过 __threadfence() 完成的。