c++ - 为什么我在 CUDA 中实现求和减少时会得到错误的结果？

标签 c++ cuda sum reduction

我正在编写一个使用 CUDA C++ API 实现的 vector 缩减算法的教程，我很挣扎，因为我真的不明白我做错了什么，因为结果是 (device: 4386.000000 host: 260795.000000)

我使用的代码如下(问题大小固定为 512)。

编辑:不幸的是问题尚未解决，我仍然得到相同的结果。我已经更新了代码，提供了完整的代码。目标是相同的，即对包含 512 个元素的 float 组的所有元素求和。 

    #define NUM_ELEMENTS 512
__global__ void reduction(float *g_data, int n)
{
        __shared__ float s_data[NUM_ELEMENTS];
      int tid = threadIdx.x;
      int index = tid + blockIdx.x*blockDim.x;
      s_data[tid] = 0.0;
      if (index < n){
        s_data[tid] = g_data[index];
      }
      __syncthreads();

      for (int s = 2; s <= blockDim.x; s = s * 2){
        if ((tid%s) == 0){
          s_data[tid] += s_data[tid + s / 2];
        }
        __syncthreads();
      }

      if (tid == 0){
        g_data[blockIdx.x] = s_data[tid];
      }
}


    // includes, system
#include <cuda_runtime.h>
#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>
#include <float.h>

// includes, kernels
#include "vector_reduction_kernel.cu"

// For simplicity, just to get the idea in this MP, we're fixing the problem size to 512 elements.
#define NUM_ELEMENTS 512

////////////////////////////////////////////////////////////////////////////////
// declaration, forward
void runTest( int argc, char** argv);

float computeOnDevice(float* h_data, int array_mem_size);

extern "C" 
void computeGold( float* reference, float* idata, const unsigned int len);

////////////////////////////////////////////////////////////////////////////////
// Program main
////////////////////////////////////////////////////////////////////////////////
int main( int argc, char** argv) 
{

cudaSetDevice(0);
    runTest( argc, argv);
    return EXIT_SUCCESS;
}

////////////////////////////////////////////////////////////////////////////////
//! Run naive scan test
////////////////////////////////////////////////////////////////////////////////
void runTest( int argc, char** argv) 
{
    int num_elements = NUM_ELEMENTS;

    const unsigned int array_mem_size = sizeof( float) * num_elements;

    // Allocate host memory to store the input data
    float* h_data = (float*) malloc( array_mem_size);

    // initialize the input data on the host to be integer values
    // between 0 and 1000
    for( unsigned int i = 0; i < num_elements; ++i) 
        h_data[i] = floorf(1000*(rand()/(float)RAND_MAX));

    // Function to compute the reference solution on CPU using a C sequential version of the algorithm
    // It is written in the file "vector_reduction_gold.cpp". The Makefile compiles this file too.
    float reference = 0.0f;  
    computeGold(&reference , h_data, num_elements);

    // Function to compute the solution on GPU using a call to a CUDA kernel (see body below)
    // The kernel is written in the file "vector_reduction_kernel.cu". The Makefile also compiles this file.
    float result = computeOnDevice(h_data, num_elements);

    // We can use an epsilon of 0 since values are integral and in a range that can be exactly represented
    float epsilon = 0.0f;
    unsigned int result_regtest = (abs(result - reference) <= epsilon);
    printf( "Test %s\n", (1 == result_regtest) ? "Ok." : "No.");
    printf( "device: %f  host: %f\n", result, reference);
    // cleanup memory
    free( h_data);
}

// Function to call the CUDA kernel on the GPU.
// Take h_data from host, copies it to device, setup grid and thread 
// dimensions, excutes kernel function, and copy result of scan back
// to h_data.
// Note: float* h_data is both the input and the output of this function.
float computeOnDevice(float* h_data, int num_elements)
{
  float* d_data = NULL;
  float result;

  // Memory allocation on device side
  cudaMalloc((void**)&d_data, sizeof(float)*num_elements);

  // Copy from host memory to device memory
  cudaMemcpy((void**)&d_data, h_data, num_elements * sizeof(float), cudaMemcpyHostToDevice );

  //int threads = (num_elements/2) + num_elements%2;
  int threads = (num_elements);
  // Invoke the kernel
  reduction<<< 1 ,threads >>>(d_data,num_elements);

  // Copy from device memory back to host memory
  cudaMemcpy(&result, d_data, sizeof(float), cudaMemcpyDeviceToHost);

  cudaFree(d_data);
  cudaDeviceReset();
  return result;
}

float computeOnDevice(float* h_data, int num_elements)
    {
      float* d_data = NULL;
      float result;

      // Memory allocation on device side
      cudaMalloc((void**)&d_data, sizeof(float)*num_elements);

      // Copy from host memory to device memory
      cudaMemcpy(d_data, h_data, num_elements * sizeof(float), cudaMemcpyHostToDevice );

      int threads = (num_elements);

      // Invoke the kernel
      reduction<<< 1 ,threads >>>(d_data,num_elements);

      // Copy from device memory back to host memory
      cudaMemcpy(&result, d_data, sizeof(float), cudaMemcpyDeviceToHost);
      cudaFree(d_data);
      cudaDeviceReset();
      return result;
    }

最佳答案

您确实应该为此类问题提供完整的代码。您还应该使用 proper CUDA error checking并使用cuda-memcheck运行您的代码。您的代码中至少有 2 个错误:

  1. 我们不会像这样执行cudaMemcpy:

      cudaMemcpy((void**)&d_data, h_data, num_elements * sizeof(float), cudaMemcpyHostToDevice );

    应该是:

      cudaMemcpy(d_data, h_data, num_elements * sizeof(float), cudaMemcpyHostToDevice );

    第一个参数只是一个指针，而不是指针到指针。 cuda-memcheck 或正确的 CUDA 错误检查会将您的注意力集中在这一行上。

  2. 您没有启动足够的线程。您的内核为每个线程加载一个元素。如果问题大小为 512，则将需要 512 个线程，这:

      int threads = (num_elements/2) + num_elements%2;

    没让你明白。不确定你在那里有什么想法。但这可以修复 512 情况:

      int threads = (num_elements);

    您的缩减方法需要二的幂的线程 block 大小。

这是一个完整的测试用例，注意使用cuda-memcheck:

$ cat t27.cu
#include <stdio.h>
        #define NUM_ELEMENTS 512
    __global__ void reduction(float *g_data, int n)
    {
        __shared__ float s_data[NUM_ELEMENTS];
      int tid = threadIdx.x;
      int index = tid + blockIdx.x*blockDim.x;
      s_data[tid] = 0.0;
      if (index < n){
        s_data[tid] = g_data[index];
      }
      __syncthreads();

      for (int s = 2; s <= blockDim.x; s = s * 2){
        if ((tid%s) == 0){
          s_data[tid] += s_data[tid + s / 2];
        }
        __syncthreads();
      }

      if (tid == 0){
        g_data[blockIdx.x] = s_data[tid];
      }
    }

float computeOnDevice(float* h_data, int num_elements)
    {
      float* d_data = NULL;
      float result;

      // Memory allocation on device side
      cudaMalloc((void**)&d_data, sizeof(float)*num_elements);

      // Copy from host memory to device memory
      cudaMemcpy(d_data, h_data, num_elements * sizeof(float), cudaMemcpyHostToDevice );

      int threads = (num_elements);

      // Invoke the kernel
      reduction<<< 1 ,threads >>>(d_data,num_elements);

      // Copy from device memory back to host memory
      cudaMemcpy(&result, d_data, sizeof(float), cudaMemcpyDeviceToHost);
      cudaFree(d_data);
      cudaDeviceReset();
      return result;
    }


int main(){

   float *data = new float[NUM_ELEMENTS];
   for (int i = 0; i < NUM_ELEMENTS; i++) data[i] = 1;
   float r = computeOnDevice(data, NUM_ELEMENTS);
   printf(" result = %f\n" , r);
}
$ nvcc -arch=sm_35 -o t27 t27.cu
$ cuda-memcheck ./t27
========= CUDA-MEMCHECK
 result = 512.000000
========= ERROR SUMMARY: 0 errors

这是您现在发布的代码的修改版本(以几种新的/不同的方式被破坏)，它似乎对我来说可以正确运行:

$ cat t30.cu
    #define NUM_ELEMENTS 512
__global__ void reduction(float *g_data, int n)
{
        __shared__ float s_data[NUM_ELEMENTS];
      int tid = threadIdx.x;
      int index = tid + blockIdx.x*blockDim.x;
      s_data[tid] = 0.0;
      if (index < n){
        s_data[tid] = g_data[index];
      }
      __syncthreads();

      for (int s = 2; s <= blockDim.x; s = s * 2){
        if ((tid%s) == 0){
          s_data[tid] += s_data[tid + s / 2];
        }
        __syncthreads();
      }

      if (tid == 0){
        g_data[blockIdx.x] = s_data[tid];
      }
}


    // includes, system
#include <cuda_runtime.h>
#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>
#include <float.h>

// includes, kernels

// For simplicity, just to get the idea in this MP, we're fixing the problem size to 512 elements.
#define NUM_ELEMENTS 512

////////////////////////////////////////////////////////////////////////////////
// declaration, forward
void runTest( int argc, char** argv);

float computeOnDevice(float* h_data, int array_mem_size);

extern "C"
void computeGold( float* reference, float* idata, const unsigned int len)
{
  for (int i = 0; i<len; i++) *reference += idata[i];
};

////////////////////////////////////////////////////////////////////////////////
// Program main
////////////////////////////////////////////////////////////////////////////////
int main( int argc, char** argv)
{

cudaSetDevice(0);
    runTest( argc, argv);
    return EXIT_SUCCESS;
}

////////////////////////////////////////////////////////////////////////////////
//! Run naive scan test
////////////////////////////////////////////////////////////////////////////////
void runTest( int argc, char** argv)
{
    int num_elements = NUM_ELEMENTS;

    const unsigned int array_mem_size = sizeof( float) * num_elements;

    // Allocate host memory to store the input data
    float* h_data = (float*) malloc( array_mem_size);

    // initialize the input data on the host to be integer values
    // between 0 and 1000
    for( unsigned int i = 0; i < num_elements; ++i)
        h_data[i] = floorf(1000*(rand()/(float)RAND_MAX));

    // Function to compute the reference solution on CPU using a C sequential version of the algorithm
    // It is written in the file "vector_reduction_gold.cpp". The Makefile compiles this file too.
    float reference = 0.0f;
    computeGold(&reference , h_data, num_elements);

    // Function to compute the solution on GPU using a call to a CUDA kernel (see body below)
    // The kernel is written in the file "vector_reduction_kernel.cu". The Makefile also compiles this file.
    float result = computeOnDevice(h_data, num_elements);

    // We can use an epsilon of 0 since values are integral and in a range that can be exactly represented
    float epsilon = 0.0f;
    unsigned int result_regtest = (abs(result - reference) <= epsilon);
    printf( "Test %s\n", (1 == result_regtest) ? "CORRECTO: Coinciden los resultados de la CPU y la GPU" : "INCORRECTO: Los resultados calculados en paralelo en la GPU no coinciden con los obtenidos secuencialmente en la CPU");
    printf( "device: %f  host: %f\n", result, reference);
    // cleanup memory
    free( h_data);
}

// Function to call the CUDA kernel on the GPU.
// Take h_data from host, copies it to device, setup grid and thread
// dimensions, excutes kernel function, and copy result of scan back
// to h_data.
// Note: float* h_data is both the input and the output of this function.
#if 0
float computeOnDevice(float* h_data, int num_elements)
{
  float* d_data = NULL;
  float result;

  // Memory allocation on device side
  cudaMalloc((void**)&d_data, sizeof(float)*num_elements);

  // Copy from host memory to device memory
  cudaMemcpy((void**)&d_data, h_data, num_elements * sizeof(float), cudaMemcpyHostToDevice );

  //int threads = (num_elements/2) + num_elements%2;
  int threads = (num_elements);
  // Invoke the kernel
  reduction<<< 1 ,threads >>>(d_data,num_elements);

  // Copy from device memory back to host memory
  cudaMemcpy(&result, d_data, sizeof(float), cudaMemcpyDeviceToHost);

  cudaFree(d_data);
  cudaDeviceReset();
  return result;
}
#endif
float computeOnDevice(float* h_data, int num_elements)
    {
      float* d_data = NULL;
      float result;

      // Memory allocation on device side
      cudaError_t err = cudaMalloc((void**)&d_data, sizeof(float)*num_elements);
      if (err != cudaSuccess) {printf("CUDA error: %s\n", cudaGetErrorString(err)); exit(0);}
      // Copy from host memory to device memory
      cudaMemcpy(d_data, h_data, num_elements * sizeof(float), cudaMemcpyHostToDevice );

      int threads = (num_elements);

      // Invoke the kernel
      reduction<<< 1 ,threads >>>(d_data,num_elements);

      // Copy from device memory back to host memory
      cudaMemcpy(&result, d_data, sizeof(float), cudaMemcpyDeviceToHost);
      cudaFree(d_data);
      err = cudaGetLastError();
      if (err != cudaSuccess) {printf("CUDA error: %s\n", cudaGetErrorString(err)); exit(0);}
      cudaDeviceReset();
      return result;
    }
$ nvcc -arch=sm_35 -o t30 t30.cu
$ cuda-memcheck ./t30
========= CUDA-MEMCHECK
Test CORRECTO: Coinciden los resultados de la CPU y la GPU
device: 260795.000000  host: 260795.000000
========= ERROR SUMMARY: 0 errors
$

您仍然没有在代码中添加正确的 CUDA 错误检查，因此完全有可能存在机器设置问题。如果您仍然遇到问题，您可能需要运行我上面发布的确切代码，因为我已经在其中进行了基本的错误检查。

关于c++ - 为什么我在 CUDA 中实现求和减少时会得到错误的结果？，我们在Stack Overflow上找到一个类似的问题： https://stackoverflow.com/questions/48242903/

上一篇：c++ - 将 vector<vector<int>> 的最后一个元素移动到开头

下一篇：c++ - std::mutex 是否足以使所有先前的读写发生在同一线程中的后续读写之前？

相关文章：

cuda - 如何替换已弃用的 tex2D(texture<T, 2, cudaReadModeElementType>, float, float) [with T=float]？

javascript - React JS获取数组中数字的总和

sqlite - phonegap sqlite数据库总和查询结果

c++ - 如何将 write() 系统调用设置为 "flush"？

c++ - cpp代码在什么情况下必须编译链接成exe

caching - 英伟达 CUDA : cache L2 and multiple kernel invocations

cuda - 监控线程 block 在执行时间内如何分配给 SM？

c# - List<int> 中 int 的总和范围

c++ - 自定义快捷方式的上下文菜单

c++ - 真实光可视化(数学相关)

©2024 IT工具网  联系我们