1

我是 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;
  }

我认为结果应该是 avg[0]=640.000 avg[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 示例中有更多的函数被认为是未定义的。

4

1 回答 1

5

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

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

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

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

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

您可能希望查看 C 编程指南的原子操作部分。

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

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

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

#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() （比较和交换）。 ”。

于 2012-12-12T17:08:42.080 回答