CUDA程序优化小记（二）

CUDA程序优化小记（二）

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

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

       以前写了一段时间多线程程序，有用Windows的API写的，也有用跨平台Boost写的。其实多线程的优势在于能够提升资源利用率。比如说如果我没有其它耗时的任务，那么可以开多线程进行多个任务的操作。这种情况一般发生在初始化、大量载入资源的时候。其实GPU是一种并行处理器，它被设计为能执行简单的指令，但是并行度非常高的结构，由这些特点可以看出，GPU非常适合做并行程序。要不然怎么会有CUDA中“网格 – 块 –线程”这样的结构呢。

       但是像上一次执行单线程的做法，无疑会闲置非常多的并行计算资源，我们必须采用多线程，将人物量分组，这样效率就会得到大幅度的提升。在GPU编程中，有一个叫“掩藏”（Hide）的概念。它的意思是指线程因为访问存储器或者阻塞等其它原因造成的延迟，但由于分配的线程足够多，导致整体上看GPU仍然处于忙碌的状态。

       这里顺便介绍一下锁页存储器的概念。通常来说，我们使用new或malloc函数进行分配内存操作时，一般操作系统会以页为单位进行内存的分配，这些内存页通过分页算法能够正确地映射到内存中，为的是合理使用内存。而锁页存储器是一种不受分页影响的存储器，由于这一特性，使得读取的速度非常快，因而适合进行主机（CPU）到设备（GPU）之间的传输。CUDA中有cudaMallocHost()函数进行锁页存储的分配。

       下面就是我改进自上一版后的程序源码，使用了多线程分组缩减。具体实现方法是定义了THREAD_NUM这个宏，指明使用了多少个线程执行操作，这里指定的操作线程是256个，随后内核执行函数中的任务量进行了修改。由于这些线程是并行的，因此只需要其中的一个线程计时即可。最后调用的方法中参数也须作相应的修改：

Kernel_SquareSum<<<1,THREAD_NUM>>>( pDevIn, pDevDataSize, pDevOut, pDevElasped );

    这里使用的是1个网格，THREAD_NUM个块，每个块中各有一个线程在运行。下面是所有代码：

#include <cuda_runtime.h>#include <cctype>#include <cassert>#include <cstdio>#include <ctime>#define DATA_SIZE 1048576#define THREAD_NUM 256#ifndef nullptr#define nullptr 0#endifusing namespace std;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;}}////////////////////////在设备上运行的内核函数/////////////////////////////__global__ static void Kernel_SquareSum( int* pIn, size_t* pDataSize,int* pOut, clock_t* pElapsed ){// 开始计时clock_t startTime = clock( );for ( size_t i = 0; i < *pDataSize; ++i ){*pOut += pIn[i] * pIn[i];}*pElapsed = clock( ) - startTime;// 结束计时，返回至主程序}bool CUDA_SquareSum( int* pOut, clock_t* pElapsed,int* pIn, size_t dataSize ){assert( pIn != nullptr );assert( pOut != nullptr );int* pDevIn = nullptr;int* pDevOut = nullptr;size_t* pDevDataSize = nullptr;clock_t* pDevElasped = 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, 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**)&pDevElasped, 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<<<1, 1>>>( pDevIn, pDevDataSize, pDevOut, pDevElasped );// 5、查询内核初始化的时候是否出错cudaStatus = cudaGetLastError( );if ( cudaStatus != cudaSuccess ){fprintf( stderr, "显卡执行程序时失败！" );break;}// 6、与内核同步等待执行完毕cudaStatus = cudaDeviceSynchronize( );if ( cudaStatus != cudaSuccess ){fprintf( stderr, "在与内核同步的过程中发生问题！" );break;}// 7、获取数据cudaStatus = cudaMemcpy( pOut, pDevOut, sizeof( int ), cudaMemcpyDeviceToHost );if ( cudaStatus != cudaSuccess ){fprintf( stderr, "在将结果数据从显卡复制到宿主程序中失败！" );break;}cudaStatus = cudaMemcpy( pElapsed, pDevElasped, sizeof( clock_t ), cudaMemcpyDeviceToHost );if ( cudaStatus != cudaSuccess ){fprintf( stderr, "在将耗费用时数据从显卡复制到宿主程序中失败！" );break;}cudaFree( pDevIn );cudaFree( pDevOut );cudaFree( pDevDataSize );cudaFree( pDevElasped );return true;}cudaFree( pDevIn );cudaFree( pDevOut );cudaFree( pDevDataSize );cudaFree( pDevElasped );return false;}int main( int argc, char** argv )// 函数的主入口{int* pData = nullptr;int* pResult = nullptr;clock_t* pElapsed = nullptr;// 使用CUDA内存分配器分配host端cudaError_t cudaStatus = cudaMallocHost( &pData, DATA_SIZE * sizeof( int ) );if ( cudaStatus != cudaSuccess ){fprintf( stderr, "在主机中分配资源失败！" );return 1;}cudaStatus = cudaMallocHost( &pResult, sizeof( int ) );if ( cudaStatus != cudaSuccess ){fprintf( stderr, "在主机中分配资源失败！" );return 1;}cudaStatus = cudaMallocHost( &pElapsed, sizeof( clock_t ) );if ( cudaStatus != cudaSuccess ){fprintf( stderr, "在主机中分配资源失败！" );return 1;}GenerateData( pData, DATA_SIZE );// 通过随机数产生数据CUDA_SquareSum( pResult, pElapsed, pData, DATA_SIZE );// 执行平方和// 判断是否溢出char* pOverFlow = nullptr;if ( *pResult < 0 ) pOverFlow = "（溢出）";else pOverFlow = "";// 显示基准测试printf( "用CUDA计算平方和的结果是：%d%s\n耗费用时：%d\n",*pResult, pOverFlow, *pElapsed );cudaDeviceProp prop;if ( cudaGetDeviceProperties( &prop, 0 ) == cudaSuccess ){clock_t actualTime = *pElapsed / clock_t( prop.clockRate );printf( "实际执行时间为：%dms\n", actualTime );printf( "带宽为：%.2fMB/s\n",float( DATA_SIZE * sizeof( int ) >> 20 ) * 1000.0f / float( actualTime ) );printf( "GPU设备型号：%s\n", prop.name );}cudaFreeHost( pData );cudaFreeHost( pResult );cudaFreeHost( pElapsed );return 0;}

程序的运行结果出人意料，在我的显卡GeForce GT750M上跑出了这样的成绩：

       下面对比一下三种显卡的执行效率：

显卡

执行时间

带宽

GeForce 9500 GT

11ms

363.64MB/s

GeForce 9600M GT

11ms

363.81MB/s

GeForce GT750M

4ms

1000.0MB/s