CUDA的volatile和threadfence。

Question

CUDA的volatile和threadfence。

4

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

__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。对于原子版本（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 是一个用于实现全局屏障的函数。由于存在块间通信，因此我必须保持数据一致性。

- Xuhao Chen

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Giulio Franco · Accepted Answer

CUDA编程指南中提到：

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

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

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

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

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

此外，在同一个warp中的线程之间，volatile可作为__threadfence()的快速版本（因为warp中的线程同步执行）。