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

Question

我是cuda的新手,我有一个问题.我想对我的线程进行同步,所以我尝试使用syncthreads.问题是Visual Studio 2010说:idetifier __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;
  }

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

但不仅我的结果不同(这可能是一个溢出问题),但它们并不稳定.例如,对于三种不同的执行,结果是:

avg [0] = 3041 avg [1] = 6604

avg [0] = 3015 avg [1] = 6578

avg [0] = 3047 avg [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编程指南的原子操作部分.

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

另请注意,您发布的代码不能完整编译,有许多缺少的定义,以及代码的各种其他问题.但是,由于您要发布结果,我认为您有一些版本的工作,即使您还没有发布它.因此,我没有发现您发布的代码的每个问题.

这里的代码类似于你的代码,修复了所有各种编码问题,它似乎对我有用:

#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的设备上使用这里的部分以"但请注意,任何原子操作可以根据atomicCAS(已实施开始)(比较并交换)."