CUDA volatile 和 threadfence

Question

以下两个函数有什么区别？

__device__ inline void comparator_volatile(volatile float &A, volatile float &B, uint dir) {
float t;
if ((A > B) == dir) {
    t = A;
    A = B;
    B = t;
  }
}

__device__ inline void comparator(float &A, float &B, uint dir) {
float t;
if ((A > B) == dir) {
    t = A;
    A = B;
    B = t;
  }
  __threadfence();
}

有人可以帮助我吗？

我基于 CUDA SDK 版本在一些不同的版本中实现了 BitonicSort。
对于 ATOMIC 版本 (bitonicSortAtomic)，我尝试在 __syncblocks_atomic 中使用 __threadfence() 来保持内存一致性。但它不起作用(输出不正确)。我必须调用comparator_volatile 而不是comparator，然后我得到正确的结果。任何的想法？
BitonicSort 基准测试:

// (C) Copyright 2013, University of Illinois. All Rights Reserved
#include <stdlib.h>
#include <stdio.h>
#include "parboil.h"

#define THREADS 256
#define BLOCKS 32
#define NUM_VALS 2*THREADS*BLOCKS

__device__ volatile int mutex = 0;
__device__ inline void __syncblocks_atomic(int goal) {
    __syncthreads();
//  __threadfence();
    int tx = threadIdx.x;
    if (tx == 0) {
        atomicAdd((int *)&mutex, 1);
        while(g_mutex != goal) {}
    }
    __syncthreads();
}

__device__ inline void comparator(float &A, float &B, uint dir) {
    float t;
    if ((A > B) == dir) {
        t = A;
        A = B;
        B = t;
    }
}

__device__ inline void comparator_volatile(volatile float &A, volatile float &B, uint dir) {
    float t;
    if ((A > B) == dir) {
        t = A;
        A = B;
        B = t;
    }
}

#ifdef NAIVE
__global__ void bitonicSortNaive(float *src, int stride, int size) {
  unsigned int tid = threadIdx.x + blockDim.x * blockIdx.x;
  uint dir = (tid & (size / 2)) == 0;
  unsigned int pos = 2*tid - (tid & (stride - 1));
  comparator(src[pos], src[pos+stride], dir);
}
#endif

#ifdef ATOMIC
__global__ void bitonicSortAtomic(float *src, int length) {
  uint numBlocks = gridDim.x * gridDim.y * gridDim.z;
  uint goalVal = 0;
  uint tid = threadIdx.x + blockDim.x * blockIdx.x;
  for(uint size=2; size<=length; size<<=1) {
    for(uint stride=size>>1; stride>0; stride=stride>>1) {
      uint dir = (tid & (size / 2)) == 0;
      uint pos = 2*tid - (tid & (stride - 1));
      comparator_volatile(src[pos], src[pos+stride], dir);
      if(stride>THREADS || (stride==1 && size>=THREADS)) {
        goalVal += numBlocks;
        __syncblocks_atomic(goalVal);
      }
      else
        __syncthreads();
    } // end for stride
  } // end for size
}
#endif

int main() {
  printf("[BENCH] Bitonic Sort %d elements\n", NUM_VALS);
  printf("[BENCH] Xuhao Chen <cxh@illinois.edu>\n");
#ifdef NAIVE
  printf("[BENCH] Naive version\n");
#endif
#ifdef ATOMIC
  printf("[BENCH] Atomic Barrier\n");
#endif
  float *values = (float*) malloc( NUM_VALS * sizeof(float));
  array_init(values, NUM_VALS);
  float *dev_values;
  size_t size = NUM_VALS * sizeof(float);
  cudaMalloc((void**) &dev_values, size);
  cudaMemcpy(dev_values, values, size, cudaMemcpyHostToDevice);
  dim3 blocks(BLOCKS,1);
  dim3 threads(THREADS,1);
  cudaDeviceSynchronize();

#ifdef NAIVE
  int j, k;
  for (k = 2; k <= NUM_VALS; k <<= 1) {
    for (j=k>>1; j>0; j=j>>1) {
      bitonicSortNaive<<<blocks, threads>>>(dev_values, j, k);
    }
  }
#endif

#ifdef ATOMIC
  bitonicSortAtomic<<<blocks, threads>>>(dev_values, NUM_VALS);
#endif

  cudaDeviceSynchronize();
  cudaMemcpy(values, dev_values, size, cudaMemcpyDeviceToHost);
  cudaFree(dev_values);
  free(values);
}

__syncblocks_atomic 是一个实现全局屏障的函数。由于有 block 间通信，我必须保持数据一致性。

Answer 1

CUDA 编程指南指出:

If a variable located in global or shared memory is declared as volatile, the compiler assumes that its value can be changed or used at any time by another thread and therefore any reference to this variable compiles to an actual memory read or write instruction.

这基本上意味着当您为变量赋值时内存将被立即刷新，并且当您尝试读取其值时将直接从内存中获取(没有缓存)。

在您的第一个代码示例中，由于 A 和 B 都是 volatile 的，因此生成了 6 条实际的内存指令。每次使用 A 或 B 时进行一次读/写。好处是其他线程将能够在进行修改时更早地看到这些修改。缺点是执行速度会变慢，因为缓存将被禁用。

另一方面，在您的第二个代码示例中，GPU 被授权使用缓存来加速其执行，直到函数结束时，它被迫发出内存写入。如果 A 和 B 都已缓存，则仅发出 2 次内存写入。缺点是其他线程可能只能在栅栏之后看到更改的值。

您应该考虑的另一件事是操作不是原子的。
如果其他线程在您的函数执行时尝试访问 A 和 B，在这两种情况下，它们可能会看到函数的部分执行。在第二个代码示例中，这种情况不太可能发生，因为线程可能会使用其缓存值，并立即刷新最终值(无论如何，您不应该依赖于此)。

此外， volatile 在同一 warp 中的线程之间用作 __threadfence() 的更快版本。 (因为经纱中的线程是同步执行的)。