学步园

CUDA程序优化小记（八）

CUDA全称Computer Unified Device Architecture（计算机统一设备架构），它的引入为计算机计算速度质的提升提供了可能，从此微型计算机也能有与大型机相当计算的能力。可是不恰当地使用CUDA技术，不仅不会让应用程序获得提升，反而会比普通CPU的计算还要慢。最近我通过学习《GPGPU编程技术》这本书，深刻地体会到了这一点，并且用CUDA
Runtime应用改写书上的例子程序来体会CUDA技术给我们计算能力带来的提升。

原创文章，反对未声明的引用。原博客地址：http://blog.csdn.net/gamesdev/article/details/18867665

       CUDA虽然能够让类C语言代码执行在GPU上，但是需要注意的是，由于GPU和CPU架构上的差异，无法将C语言上所有的概念都映射到CUDA上来。一个类似的例子是在OpenGL的着色语言GLSL中虽然能够支持流程控制（flow
control），即支持while、do while和for循环，但是循环的次数必须是固定的，而不能由变量指定。这就限制了更加复杂的计算。事实上，CPU能够支持完整的算术和逻辑指令集，并且具有分支预测的能力，能保持极低的内存误访问率。可是GPU的硬件设计使其拥有上百万个标量处理器，因此每个标量处理器的设计就变得相当简单，只能支持基本的算术和逻辑指令。同时GPU的计算要求并行，过于复杂的计算指令会使处于同一warp的线程造成潜在的阻塞，导致串行化（sequentialize），这样不利于计算效率的提升，总之了解GPU的构造和硬件计算原理对于提升CUDA程序的效率有好处。

       和CPU执行流程一样，GPU程序也会经历取指令、取数据、执行和保存数据结果这几步骤。上一篇文章讲到了有可能影响计算效率的步骤之一——取数据（data
fetches），简而言之，在使用共享存储器的时候要适当地避免bank冲突。不过在执行的过程中也会影响计算的效率，虽然这种影响相比上面并不那么显著，但是也不能完全忽略。

一般来说，使用流程控制产生的指令比顺序执行所产生的指令要多，如果尝试将两个执行效果等价但一个使用了流程控制另一个只是顺序执行的程序代码所生成的汇编代码相比较，那么结果表明使用流程控制控制所产生的汇编代码会更多。此外，过多地使用流程控制将会产生上述对warp内的线程产生阻塞，我们应当适当地规避。于是我们这次将内核代码中的流程控制改写一下，希望能够产生更少的指令以及更高的数据带宽。

       下面是修改后的程序代码：

////////////////////////在设备上运行的内核函数/////////////////////////////
__global__ static voidKernel_SquareSum( int* pIn, size_t* pDataSize,
                                   int*pOut, clock_t* pTime )
{
    // 声明一个动态分配的共享存储器
    extern __shared__ int sharedData[];
 
    const size_t computeSize =*pDataSize / THREAD_NUM;
    const size_t tID = size_t(threadIdx.x );// 线程
    const size_t bID = size_t(blockIdx.x );// 块
 
    int offset = 1;      // 记录每轮增倍的步距
 
    // 开始计时
    if ( tID == 0 ) pTime[bID] =clock( );// 选择任意一个线程进行计时
 
    // 执行计算
    for ( size_t i = bID * THREAD_NUM+ tID;
       i < DATA_SIZE;
       i += BLOCK_NUM * THREAD_NUM )
    {
       sharedData[tID] += pIn[i] * pIn[i];
    }
 
    // 同步一个块中的其它线程
    __syncthreads( );
    if ( tID < 128 )sharedData[tID] += sharedData[tID + 128];
    __syncthreads( );
    if ( tID < 64 )sharedData[tID] += sharedData[tID + 64];
    __syncthreads( );
    if ( tID < 32 )sharedData[tID] += sharedData[tID + 32];
    //__syncthreads( );
    if ( tID < 16 )sharedData[tID] += sharedData[tID + 16];
    //__syncthreads( );
    if ( tID < 8 ) sharedData[tID]+= sharedData[tID + 8];
    //__syncthreads( );
    if ( tID < 4 ) sharedData[tID]+= sharedData[tID + 4];
    //__syncthreads( );
    if ( tID < 2 ) sharedData[tID]+= sharedData[tID + 2];
    //__syncthreads( );
    if ( tID < 1 ) sharedData[tID]+= sharedData[tID + 1];
 
    if ( tID == 0 )// 如果线程ID为，那么计算结果，并记录时钟
    {
       pOut[bID] = sharedData[0];
       pTime[bID + BLOCK_NUM] = clock( );
    }
}

从上面的代码中我们可以看出我们将上次的for循环语句展开成使用若干语句顺序执行。注意我将后面的内核线程同步语句__syncthreads()的调用注释掉了，为什么会注释掉呢？因为每一个warp包含了32个线程，而warp是线程同时执行的最大单位，因此我们确定在一个warp中的线程没有必要进行同步。所以后面的同步语句自然而然地注释掉了。

       下面是程序的运行结果：

       下面是三种显卡执行的数据对比：

       可见带宽比上一版又提高了一点点。

       下面是程序的所有代码：

#include <cuda_runtime.h>
#include <cctype>
#include <cassert>
#include <cstdio>
#include <ctime>
 
#define DATA_SIZE 1048576
#define BLOCK_NUM 32
#define THREAD_NUM 256
#ifndef nullptr
#define nullptr 0
#endif
 
using namespace std;
 
////////////////////////在设备上运行的内核函数/////////////////////////////
__global__ static voidKernel_SquareSum( int* pIn, size_t* pDataSize,
                                   int*pOut, clock_t* pTime )
{
    // 声明一个动态分配的共享存储器
    extern __shared__ int sharedData[];
 
    const size_t computeSize =*pDataSize / THREAD_NUM;
    const size_t tID = size_t(threadIdx.x );// 线程
    const size_t bID = size_t(blockIdx.x );// 块
 
    int offset = 1;      // 记录每轮增倍的步距
 
    // 开始计时
    if ( tID == 0 ) pTime[bID] =clock( );// 选择任意一个线程进行计时
 
    // 执行计算
    for ( size_t i = bID * THREAD_NUM+ tID;
       i < DATA_SIZE;
       i += BLOCK_NUM * THREAD_NUM )
    {
       sharedData[tID] += pIn[i] * pIn[i];
    }
 
    // 同步一个块中的其它线程
    __syncthreads( );
    if ( tID < 128 )sharedData[tID] += sharedData[tID + 128];
    __syncthreads( );
    if ( tID < 64 )sharedData[tID] += sharedData[tID + 64];
    __syncthreads( );
    if ( tID < 32 )sharedData[tID] += sharedData[tID + 32];
    //__syncthreads( );
    if ( tID < 16 )sharedData[tID] += sharedData[tID + 16];
    //__syncthreads( );
    if ( tID < 8 ) sharedData[tID]+= sharedData[tID + 8];
    //__syncthreads( );
    if ( tID < 4 ) sharedData[tID]+= sharedData[tID + 4];
    //__syncthreads( );
    if ( tID < 2 ) sharedData[tID]+= sharedData[tID + 2];
    //__syncthreads( );
    if ( tID < 1 ) sharedData[tID]+= sharedData[tID + 1];
 
    if ( tID == 0 )// 如果线程ID为，那么计算结果，并记录时钟
    {
       pOut[bID] = sharedData[0];
       pTime[bID + BLOCK_NUM] = clock( );
    }
}
 
bool CUDA_SquareSum( int* pOut,clock_t* pTime,
                  int* pIn, size_tdataSize )
{
    assert( pIn != nullptr );
    assert( pOut != nullptr );
 
    int* pDevIn = nullptr;
    int* pDevOut = nullptr;
    size_t* pDevDataSize = nullptr;
    clock_t* pDevTime = nullptr;
 
    // 1、设置设备
    cudaError_t cudaStatus = cudaSetDevice( 0 );// 只要机器安装了英伟达显卡，那么会调用成功
    if ( cudaStatus != cudaSuccess )
    {
       fprintf( stderr, "调用cudaSetDevice()函数失败！" );
       return false;
    }
 
    switch ( true)
    {
    default:
       // 2、分配显存空间
       cudaStatus = cudaMalloc( (void**)&pDevIn,dataSize * sizeof( int) );
       if ( cudaStatus != cudaSuccess)
       {
           fprintf( stderr, "调用cudaMalloc()函数初始化显卡中数组时失败！" );
           break;
       }
 
       cudaStatus = cudaMalloc( (void**)&pDevOut,BLOCK_NUM * sizeof( int) );
       if ( cudaStatus != cudaSuccess)
       {
           fprintf( stderr, "调用cudaMalloc()函数初始化显卡中返回值时失败！" );
           break;
       }
 
       cudaStatus = cudaMalloc( (void**)&pDevDataSize,sizeof( size_t ) );
       if ( cudaStatus != cudaSuccess)
       {
           fprintf( stderr, "调用cudaMalloc()函数初始化显卡中数据大小时失败！" );
           break;
       }
 
       cudaStatus = cudaMalloc( (void**)&pDevTime,BLOCK_NUM * 2 * sizeof( clock_t ) );
       if ( cudaStatus != cudaSuccess)
       {
           fprintf( stderr, "调用cudaMalloc()函数初始化显卡中耗费用时变量失败！" );
           break;
       }
 
       // 3、将宿主程序数据复制到显存中
       cudaStatus = cudaMemcpy( pDevIn, pIn, dataSize * sizeof( int ),cudaMemcpyHostToDevice );
       if ( cudaStatus != cudaSuccess)
       {
           fprintf( stderr, "调用cudaMemcpy()函数初始化宿主程序数据数组到显卡时失败！" );
           break;
       }
 
       cudaStatus = cudaMemcpy( pDevDataSize, &dataSize, sizeof( size_t ), cudaMemcpyHostToDevice );
       if ( cudaStatus != cudaSuccess)
       {
           fprintf( stderr, "调用cudaMemcpy()函数初始化宿主程序数据大小到显卡时失败！" );
           break;
       }
 
       // 4、执行程序，宿主程序等待显卡执行完毕
       Kernel_SquareSum<<<BLOCK_NUM, THREAD_NUM, THREAD_NUM* sizeof( int)>>>
           ( pDevIn, pDevDataSize, pDevOut, pDevTime );
 
       // 5、查询内核初始化的时候是否出错
       cudaStatus = cudaGetLastError( );
       if ( cudaStatus != cudaSuccess)
       {
           fprintf( stderr, "显卡执行程序时失败！" );
           break;
       }
 
       // 6、与内核同步等待执行完毕
       cudaStatus = cudaDeviceSynchronize( );
       if ( cudaStatus != cudaSuccess)
       {
           fprintf( stderr, "在与内核同步的过程中发生问题！" );
           break;
       }
 
       // 7、获取数据
       cudaStatus = cudaMemcpy( pOut, pDevOut, BLOCK_NUM * sizeof( int ),cudaMemcpyDeviceToHost );
       if ( cudaStatus != cudaSuccess)
       {
           fprintf( stderr, "在将结果数据从显卡复制到宿主程序中失败！" );
           break;
       }
 
       cudaStatus = cudaMemcpy( pTime, pDevTime, BLOCK_NUM * 2 * sizeof( clock_t ), cudaMemcpyDeviceToHost );
       if ( cudaStatus != cudaSuccess)
       {
           fprintf( stderr, "在将耗费用时数据从显卡复制到宿主程序中失败！" );
           break;
       }
 
       // 8、释放空间
       cudaFree( pDevIn );
       cudaFree( pDevOut );
       cudaFree( pDevDataSize );
       cudaFree( pDevTime );
       return true;
    }
 
    // 8、释放空间
    cudaFree( pDevIn );
    cudaFree( pDevOut );
    cudaFree( pDevDataSize );
    cudaFree( pDevTime );
    return false;
}
 
void GenerateData( int* pData,size_t dataSize )// 产生数据
{
    assert( pData != nullptr );
    for ( size_t i = 0; i <dataSize; i++ )
    {
       srand( i + 3 );
       pData[i] = rand( ) % 100;
    }
}
 
int main( int argc, char** argv )// 函数的主入口
{
    int* pData = nullptr;
    int* pResult = nullptr;
    clock_t* pTime = nullptr;
 
    // 使用CUDA内存分配器分配host端
    cudaError_t cudaStatus = cudaMallocHost( &pData, DATA_SIZE * sizeof( int ) );
    if ( cudaStatus != cudaSuccess )
    {
       fprintf( stderr, "在主机中分配资源失败！" );
       return 1;
    }
 
    cudaStatus = cudaMallocHost( &pResult, BLOCK_NUM * sizeof( int ) );
    if ( cudaStatus != cudaSuccess )
    {
       fprintf( stderr, "在主机中分配资源失败！" );
       return 1;
    }
 
    cudaStatus = cudaMallocHost( &pTime, BLOCK_NUM * 2 * sizeof( clock_t ) );
    if ( cudaStatus != cudaSuccess )
    {
       fprintf( stderr, "在主机中分配资源失败！" );
       return 1;
    }
 
    GenerateData( pData, DATA_SIZE );// 通过随机数产生数据
    CUDA_SquareSum( pResult, pTime, pData, DATA_SIZE );// 执行平方和
 
    // 在CPU中将结果组合起来
    int totalResult;
    for ( inti = 0; i < BLOCK_NUM; ++i )
    {
       totalResult += pResult[i];
    }
 
    // 计算执行的时间
    clock_t startTime = pTime[0];
    clock_t endTime = pTime[BLOCK_NUM];
    for ( inti = 0; i < BLOCK_NUM; ++i )
    {
       if ( startTime > pTime[i] )startTime = pTime[i];
       if ( endTime < pTime[i +BLOCK_NUM] ) endTime = pTime[i + BLOCK_NUM];
    }
    clock_t elapsed = endTime - startTime;
 
 
    // 判断是否溢出
    char* pOverFlow = nullptr;
    if ( totalResult < 0 )pOverFlow = "（溢出）";
    else pOverFlow = "";
 
    // 显示基准测试
    printf( "用CUDA计算平方和的结果是：%d%s\n耗费用时：%d\n",
       totalResult, pOverFlow, elapsed );
 
    cudaDeviceProp prop;
    if ( cudaGetDeviceProperties(&prop, 0 ) == cudaSuccess )
    {
       float actualTime = float( elapsed ) / float(prop.clockRate );
       printf( "实际执行时间为：%.2fms\n", actualTime );
       printf( "带宽为：%.2fMB/s\n",
           float( DATA_SIZE * sizeof( int ) >>20 ) * 1000.0f / actualTime );
       printf( "GPU设备型号：%s\n", prop.name );
    }
 
    cudaFreeHost( pData );
    cudaFreeHost( pResult );
    cudaFreeHost( pTime );
 
    return 0;
}