Cuda __syncthreads 未定义。没有它-> 随机结果

Question

我是 cuda 新手，但遇到了问题。我想对我的线程进行同步，所以我尝试使用同步线程。问题是 Visual Studio 2010 说：标识符 __syncthreads() 未定义...顺便说一句，我正在使用 cuda 4.2。所以我决定改用 cudaDeviceSynchronize() 并从主机调用它。我的代码类似于上面的代码（我只将重要部分发送给您）：

__global__ void sum( float avg[]){
  avg[0]+=1;
  avg[1]+=2;
}
int main(){
  float avg[2];
  float *devAvg;
  cudaError_t cudaStatus;
  size_t size=sizeof(unsigned char)*2;
  cudaStatus = cudaMalloc((void**)&devAvg, size2);
  if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMalloc 2 failed!");
    return -1;
  }
  avg[0]=0;
  avg[1]=0;
  cudaStatus = cudaMemcpy(devAvg,avg, size, cudaMemcpyHostToDevice);
  if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMemcpy failed!");
    return -1;
  }
  dim3 nblocks(40,40);
  dim3 nthreads(20,20);
  sum<<<nblocks,nthreads,msBytes>>>(devAvg);
  cudaStatus = cudaDeviceSynchronize();
  if (cudaStatus != cudaSuccess) {
      fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
  }

  cudaStatus = cudaMemcpy(avg,devAvg,size,cudaMemcpyDeviceToHost);
  if (cudaStatus != cudaSuccess) {
      fprintf(stderr, "cudaMemcpy Device to Host failed!");
      return -1;}
  cout<<"avg[0]="avg[0]<<" avg[1]="<<avg[1]<<endl;
  cudaFree devAvg;
  return 0;
  }

我认为结果应该是 平均[0]=640.000 平均[1]=1.280.000

但不仅我的结果不同（这可能是溢出问题），而且它们并不稳定。例如，对于三种不同的执行，结果是：

平均[0]=3041 平均[1]=6604

平均[0]=3015 平均[1]=6578

平均[0]=3047 平均[1]=6600

那么我在这里做错了什么？是同步问题吗？为什么我不能使用 __syncthreads() 还是竞争条件的问题？

对于 __syncthreads() 问题，我编写的任何代码都附带它。即使是最简单的：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <Windows.h>


// Kernel that executes on the CUDA device
__global__ void square_array(float *a, int N)
{
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  if (idx<N) a[idx] = a[idx] * a[idx];
  __syncthreads();
}

// main routine that executes on the host
int main(void)
{
  float *a_h, *a_d;  // Pointer to host & device arrays
  const int N = 10;  // Number of elements in arrays
  size_t size = N * sizeof(float);
  a_h = (float *)malloc(size);        // Allocate array on host
  cudaMalloc((void **) &a_d, size);   // Allocate array on device
  // Initialize host array and copy it to CUDA device
  for (int i=0; i<N; i++) a_h[i] = (float)i;
  cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice);
  // Do calculation on device:
  int block_size = 4;
  int n_blocks = N/block_size + (N%block_size == 0 ? 0:1);
  square_array <<< n_blocks, block_size >>> (a_d, N);
  // Retrieve result from device and store it in host array
  cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost);
  // Print results
  for (int i=0; i<N; i++) printf("%d %f\n", i, a_h[i]);
  // Cleanup
  free(a_h); cudaFree(a_d);
      return 0;
}

它是这样说的：错误：标识符“__syncthreads()”未定义

有趣的是，即使使用 4.2 CUDA SDK 附带的示例代码，也会发生同样的事情......也许是更普遍的错误，因为 SDK 示例中有更多被认为未定义的函数。

Answer 1

您所有的线程块都在写入相同的两个位置。使其正常工作的唯一方法是使用原子操作。否则，线程读取该位置、添加到该位置并将结果“同时”写回该位置的结果是未定义的。

如果你重写你的内核如下：

__global__ void sum( float avg[]){
   atomicAdd(&(avg[0]),1);
   atomicAdd(&(avg[1]),2);
}

它应该可以解决您看到的问题。

要回答有关 __syncthreads() 的问题，我需要查看导致编译器错误的确切代码。如果你发布，我会更新我的答案。在这个内核中插入 __syncthreads() 调用应该没有问题，尽管它不会解决您看到的问题。

您可能希望查看 C 编程指南的 atomic operations 部分。

请注意，使用原子通常会导致代码运行速度变慢，因此应谨慎使用。但是，对于这个学习练习，它应该为您解决问题。

还请注意，您发布的代码编译不干净，有许多缺少定义，以及您的代码的各种其他问题。但是由于您要发布结果，我假设您有这个工作的某个版本，即使您没有发布它。因此，我没有发现您发布的代码的所有问题。

这里的代码与您的代码相似，修复了所有各种编码问题，它似乎对我有用：

#include <stdio.h>
#include <iostream>

#define msBytes 0

__global__ void sum( float avg[]){
  atomicAdd(&(avg[0]),1);
  atomicAdd(&(avg[1]),2);
}
int main(){
  float avg[2];
  float *devAvg;
  cudaError_t cudaStatus;
  size_t size=sizeof(float)*2;
  cudaStatus = cudaMalloc((void**)&devAvg, size);
  if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMalloc 2 failed!");
    return -1;
  }
  avg[0]=0;
  avg[1]=0;
  cudaStatus = cudaMemcpy(devAvg,avg, size, cudaMemcpyHostToDevice);
  if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMemcpy failed!");
    return -1;
  }
  dim3 nblocks(40,40);
  dim3 nthreads(20,20);
  sum<<<nblocks,nthreads,msBytes>>>(devAvg);
  cudaStatus = cudaDeviceSynchronize();
  if (cudaStatus != cudaSuccess) {
      fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
  }

  cudaStatus = cudaMemcpy(avg,devAvg,size,cudaMemcpyDeviceToHost);
  if (cudaStatus != cudaSuccess) {
      fprintf(stderr, "cudaMemcpy Device to Host failed!");
      return -1;}
  std::cout<<"avg[0]="<<avg[0]<<" avg[1]="<<avg[1]<<std::endl;
  cudaFree(devAvg);
  return 0;
  }

我运行它时得到以下输出：

avg[0]=640000 avg[1]=1.28e+06

还请注意，要使atomicAdd 在float 上可用，必须具有计算能力为 2.0 或更好的设备（并通过编译器开关，例如 -arch=sm_20 以针对该设备进行编译）。如果您有较早的设备（计算能力 1.x），那么您可以创建一个类似的程序，将 avg[] 定义为 int 而不是 float。或者，如果您愿意，您可以创建自己的 atomicAdd __ device__ 函数，该函数可在 cc 1.x 设备上使用，如建议 here 在以“但请注意，任何原子操作都可以基于 atomicCAS() (比较和交换）。”。