CUDA程序优化小记（八） - 推酷
CUDA程序优化小记（八）
CUDA程序优化小记（八）
Computer Unified Device Architecture
（计算机统一设备架构），它的引入为计算机计算速度质的提升提供了可能，从此微型计算机也能有与大型机相当计算的能力。可是不恰当地使用
技术，不仅不会让应用程序获得提升，反而会比普通
的计算还要慢。最近我通过学习《
编程技术》这本书，深刻地体会到了这一点，并且用
CUDA Runtime
应用改写书上的例子程序来体会
技术给我们计算能力带来的提升。
原创文章，反对未声明的引用。原博客地址：
http://blog.csdn.net/gamesdev/article/details/
&&&&&& CUDA
虽然能够让类
语言代码执行在
上，但是需要注意的是，由于
架构上的差异，无法将
语言上所有的概念都映射到
上来。一个类似的例子是在
的着色语言
中虽然能够支持流程控制（
flow control
），即支持
循环，但是循环的次数必须是固定的，而不能由变量指定。这就限制了更加复杂的计算。事实上，
能够支持完整的算术和逻辑指令集，并且具有分支预测的能力，能保持极低的内存误访问率。可是
的硬件设计使其拥有上百万个标量处理器，因此每个标量处理器的设计就变得相当简单，只能支持基本的算术和逻辑指令。同时
的计算要求并行，过于复杂的计算指令会使处于同一
的线程造成潜在的阻塞，导致串行化（
sequentialize
），这样不利于计算效率的提升，总之了解
的构造和硬件计算原理对于提升
程序的效率有好处。
执行流程一样，
程序也会经历取指令、取数据、执行和保存数据结果这几步骤。上一篇文章讲到了有可能影响计算效率的步骤之一——取数据（
data fetches
），简而言之，在使用共享存储器的时候要适当地避免
冲突。不过在执行的过程中也会影响计算的效率，虽然这种影响相比上面并不那么显著，但是也不能完全忽略。
一般来说，使用流程控制产生的指令比顺序执行所产生的指令要多，如果尝试将两个执行效果等价但一个使用了流程控制另一个只是顺序执行的程序代码所生成的汇编代码相比较，那么结果表明使用流程控制控制所产生的汇编代码会更多。此外，过多地使用流程控制将会产生上述对
内的线程产生阻塞，我们应当适当地规避。于是我们这次将内核代码中的流程控制改写一下，希望能够产生更少的指令以及更高的数据带宽。
下面是修改后的程序代码：
////////////////////////在设备上运行的内核函数/////////////////////////////
__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( );
从上面的代码中我们可以看出我们将上次的
循环语句展开成使用若干语句顺序执行。注意我将后面的内核线程同步语句
__syncthreads()
的调用注释掉了，为什么会注释掉呢？因为每一个
个线程，而
是线程同时执行的最大单位，因此我们确定在一个
中的线程没有必要进行同步。所以后面的同步语句自然而然地注释掉了。
下面是程序的运行结果：
下面是三种显卡执行的数据对比：
GeForce 9500 GT
7488.15MB/s
GeForce 9600M GT
8647.8MB/s
GeForce GT750M
52432.53MB/s
可见带宽比上一版又提高了一点点。
下面是程序的所有代码：
#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
////////////////////////在设备上运行的内核函数/////////////////////////////
__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 =
int* pDevOut =
size_t* pDevDataSize =
clock_t* pDevTime =
// 1、设置设备
cudaError_t cudaStatus = cudaSetDevice( 0 );// 只要机器安装了英伟达显卡，那么会调用成功
if ( cudaStatus != cudaSuccess )
fprintf( stderr, &调用cudaSetDevice()函数失败！& );
switch ( true)
// 2、分配显存空间
cudaStatus = cudaMalloc( (void**)&pDevIn,dataSize * sizeof( int) );
if ( cudaStatus != cudaSuccess)
fprintf( stderr, &调用cudaMalloc()函数初始化显卡中数组时失败！& );
cudaStatus = cudaMalloc( (void**)&pDevOut,BLOCK_NUM * sizeof( int) );
if ( cudaStatus != cudaSuccess)
fprintf( stderr, &调用cudaMalloc()函数初始化显卡中返回值时失败！& );
cudaStatus = cudaMalloc( (void**)&pDevDataSize,sizeof( size_t ) );
if ( cudaStatus != cudaSuccess)
fprintf( stderr, &调用cudaMalloc()函数初始化显卡中数据大小时失败！& );
cudaStatus = cudaMalloc( (void**)&pDevTime,BLOCK_NUM * 2 * sizeof( clock_t ) );
if ( cudaStatus != cudaSuccess)
fprintf( stderr, &调用cudaMalloc()函数初始化显卡中耗费用时变量失败！& );
// 3、将宿主程序数据复制到显存中
cudaStatus = cudaMemcpy( pDevIn, pIn, dataSize * sizeof( int ),cudaMemcpyHostToDevice );
if ( cudaStatus != cudaSuccess)
fprintf( stderr, &调用cudaMemcpy()函数初始化宿主程序数据数组到显卡时失败！& );
cudaStatus = cudaMemcpy( pDevDataSize, &dataSize, sizeof( size_t ), cudaMemcpyHostToDevice );
if ( cudaStatus != cudaSuccess)
fprintf( stderr, &调用cudaMemcpy()函数初始化宿主程序数据大小到显卡时失败！& );
// 4、执行程序，宿主程序等待显卡执行完毕
Kernel_SquareSum&&&BLOCK_NUM, THREAD_NUM, THREAD_NUM* sizeof( int)&&&
( pDevIn, pDevDataSize, pDevOut, pDevTime );
// 5、查询内核初始化的时候是否出错
cudaStatus = cudaGetLastError( );
if ( cudaStatus != cudaSuccess)
fprintf( stderr, &显卡执行程序时失败！& );
// 6、与内核同步等待执行完毕
cudaStatus = cudaDeviceSynchronize( );
if ( cudaStatus != cudaSuccess)
fprintf( stderr, &在与内核同步的过程中发生问题！& );
// 7、获取数据
cudaStatus = cudaMemcpy( pOut, pDevOut, BLOCK_NUM * sizeof( int ),cudaMemcpyDeviceToHost );
if ( cudaStatus != cudaSuccess)
fprintf( stderr, &在将结果数据从显卡复制到宿主程序中失败！& );
cudaStatus = cudaMemcpy( pTime, pDevTime, BLOCK_NUM * 2 * sizeof( clock_t ), cudaMemcpyDeviceToHost );
if ( cudaStatus != cudaSuccess)
fprintf( stderr, &在将耗费用时数据从显卡复制到宿主程序中失败！& );
// 8、释放空间
cudaFree( pDevIn );
cudaFree( pDevOut );
cudaFree( pDevDataSize );
cudaFree( pDevTime );
// 8、释放空间
cudaFree( pDevIn );
cudaFree( pDevOut );
cudaFree( pDevDataSize );
cudaFree( pDevTime );
void GenerateData( int* pData,size_t dataSize )// 产生数据
assert( pData != nullptr );
for ( size_t i = 0; i &dataS i++ )
srand( i + 3 );
pData[i] = rand( ) % 100;
int main( int argc, char** argv )// 函数的主入口
int* pData =
int* pResult =
clock_t* pTime =
// 使用CUDA内存分配器分配host端
cudaError_t cudaStatus = cudaMallocHost( &pData, DATA_SIZE * sizeof( int ) );
if ( cudaStatus != cudaSuccess )
fprintf( stderr, &在主机中分配资源失败！& );
cudaStatus = cudaMallocHost( &pResult, BLOCK_NUM * sizeof( int ) );
if ( cudaStatus != cudaSuccess )
fprintf( stderr, &在主机中分配资源失败！& );
cudaStatus = cudaMallocHost( &pTime, BLOCK_NUM * 2 * sizeof( clock_t ) );
if ( cudaStatus != cudaSuccess )
fprintf( stderr, &在主机中分配资源失败！& );
GenerateData( pData, DATA_SIZE );// 通过随机数产生数据
CUDA_SquareSum( pResult, pTime, pData, DATA_SIZE );// 执行平方和
// 在CPU中将结果组合起来
int totalR
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 - startT
// 判断是否溢出
char* pOverFlow =
if ( totalResult & 0 )pOverFlow = &（溢出）&;
else pOverFlow = &&;
// 显示基准测试
printf( &用CUDA计算平方和的结果是：%d%s\n耗费用时：%d\n&,
totalResult, pOverFlow, elapsed );
cudaDeviceP
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 );
已发表评论数()
请填写推刊名
描述不能大于100个字符!
权限设置： 公开
仅自己可见
正文不准确
标题不准确
排版有问题
主题不准确
没有分页内容
图片无法显示
视频无法显示
与原文不一致以下是两个连续的CUDA核函数衔接的一种思路：要完成的功能：1. 向量的计算computer（暂时以两向量求和为例）；2. 对结果向量求和SUM。思路1：写1个计算的内核函数，中间结果保留，求和函数调用cublas。#include &cuda_runtime.h&#include &device_launch_parameters.h&#include &cublas_v2.h&#include &stdio.h&#include &stdlib.h&#pragma comment(lib, &cublas.lib&)cudaError_t addWithCuda(const int *a, const int *b, unsigned int size, float& sum);__global__ void addKernel(const int *a, const int *b, float *c){
int i = threadIdx.x;
c[i] = a[i] + b[i];}int main(){
const int arraySize = 5;
const int a[arraySize] = { 1, 2, 3, 4, 5 };
const int b[arraySize] = { 10, 20, 30, 40, 50 };
// Add vectors in parallel. float sum = 0.0;
cudaError_t cudaStatus = addWithCuda(a, b, arraySize, sum);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, &addWithCuda failed!&);
printf(&sum: %f/n&, sum);
// cudaDeviceReset must be called before exiting in order for profiling and
// tracing tools such as Nsight and Visual Profiler to show complete traces.
cudaStatus = cudaDeviceReset();
if (cudaStatus != cudaSuccess) {
fprintf(stderr, &cudaDeviceReset failed!&);
return 0;}// Helper function for using CUDA to add vectors in parallel.cudaError_t addWithCuda(const int *a, const int *b, unsigned int size, float& sum){
int *dev_a = 0;
int *dev_b = 0;
float *dev_c = 0;
cudaError_t cudaS
// Choose which GPU to run on, change this on a multi-GPU system.
cudaStatus = cudaSetDevice(0);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, &cudaSetDevice failed!
Do you have a CUDA-capable GPU installed?&);
// Allocate GPU buffers for three vectors (two input, one output)
cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(float));
if (cudaStatus != cudaSuccess) {
fprintf(stderr, &cudaMalloc failed!&);
cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
fprintf(stderr, &cudaMalloc failed!&);
cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
fprintf(stderr, &cudaMalloc failed!&);
// Copy input vectors from host memory to GPU buffers.
cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, &cudaMemcpy failed!&);
cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, &cudaMemcpy failed!&);
// Launch a kernel on the GPU with one thread for each element. addKernel && &1, size && &(dev_a, dev_b, dev_c);
// Check for any errors launching the kernel
cudaStatus = cudaGetLastError();
if (cudaStatus != cudaSuccess) {
fprintf(stderr, &addKernel launch failed: %s/n&, cudaGetErrorString(cudaStatus));
// cudaDeviceSynchronize waits for the kernel to finish, and returns
// any errors encountered during the launch.
cudaStatus = cudaDeviceSynchronize();
if (cudaStatus != cudaSuccess) {
fprintf(stderr, &cudaDeviceSynchronize returned error code %d after launching addKernel!/n&, cudaStatus);
} cublasHandle_ cublasStatus_
stat = cublasCreate(&handle); if (stat != CUBLAS_STATUS_SUCCESS) {
printf(&CUBLAS initialization failed/n&);
exit(-1); }
cublasSasum(handle, size, dev_c, 1, &sum); if (stat != CUBLAS_STATUS_SUCCESS) {
printf(&data execution failed&);
cublasDestroy(handle);
exit(-1); }
cublasDestroy(handle);Error:
cudaFree(dev_c);
cudaFree(dev_a);
cudaFree(dev_b);
return cudaS}思路2：写两个内核函数，1个调用另外一个（待求证）。__device__ void sum(const float *c, float* sum){}__global__ void calculate(const int *a, const int *b, float *c, float *sum){ //calculate.... int i = threadIdx.x; c[i] = a[i] + b[i]; __syncthreads(); sum(c, sum)}思路3：写两个内核函数，顺序执行，即自己动手实现sum函数。__global__ void calculate(const int *a, const int *b, float *c){ //calculate.... int i = threadIdx.x; c[i] = a[i] + b[i];}__global__ void sum(const float *c, float* sum){}推荐思路1，因为尽量调用库函数（不自己手写），开发效率会提高。
最新教程周点击榜
微信扫一扫CUDA编程接口:异步并发执行的概念和API
　作者: yyfn风辰　编辑:
　　【IT168技术】本文主要描述了支持系统中不同层次的异步并发执行的概念和API。　　相关阅读：　　　　　　　　　　　　1.主机和设备间异步执行　　为了易于使用主机和设备间的异步执行，一些函数是异步的：在设备完全完成任务前，控制已经返回给主机线程了。它们是： 内核发射; 设备间数据拷贝函数; 主机和设备内拷贝小于64KB的存储器块时; 存储器拷贝函数中带有Async后缀的; 设置设备存储器的函数调用。　　程序员可通过将CUDA_LAUNCH_BLOCKING环境变量设置为1来全局禁用所有运行在系统上的应用的异步内核发射。提供这个特性只是为了调试，永远不能作为使软件产品运行得可靠的方式。 当应用通过CUDA调试器或CUDA profiler(cuda-gdb, CUDA Visual Profiler, Parallel Nsight)运行时，所有的内核发射都是同步的。　　2.数据传输和内核执行重叠　　一些计算能力1.1或更高的设备可在内核执行时，在分页锁定存储器和设备存储器之间拷贝数据。应用可以通过检查asyncEngineCount 设备属性查询这种能力，如果其大于0，说明设备支持数据传输和内核执行重叠。这种能力目前只支持不涉及CUDA数组和使用cudaMallocPitch()分配的二维数组的存储器拷贝( 见前文，可阅读&相关阅读&中的文章)。　　3. 并发内核执行　　一些计算能力2.x的设备可并发执行多个内核。应用可以检查concurrentKernels属性以查询这种能力)(后续文章将介绍)，如果等于1，说明支持。 设备最大可并发执行的内核数目是16。 来自不同CUDA上下文的内核不能并发执行。 使用了许多纹理或大量本地存储器的内核和其它内核并发执行的可能性比较小。　　4. 并发数据传输　　在计算能力2.x的设备上，从主机分页锁定存储器复制数据到设备存储器和从设备存储器复制数据到主机分页锁定存储器，这两个操作可并发执行。 应用可以通过检查asyncEngineCount 属性查询这种能力，如果等于2，说明支持。　　5. 流　　应用通过流管理并发。流是一系列顺序执行的命令(可能是不同的主机线程发射)。另外，流之间相对无序的或并发的执行它们的命令;这种行为是没有保证的，而且不能作为正确性的的保证(如内核间的通信没有定义)。　　①创建和销毁　　可以通过创建流对象来定义流，且可指定它作为一系列内核发射和设备主机间存储器拷贝的流参数。下面的代码创建了两个流且在分页锁定存储器中分配了一个名为hostPtr的浮点数组。cudaStream_t stream[<span style="color: #];for (int i = <span style="color: #; i & <span style="color: #; ++i)cudaStreamCreate(&stream[i]); float* hostPcudaMallocHost((void**)&hostPtr, <span style="color: # * size);&　　下面的代码定义的每个流是一个由一次主机到设备的传输，一次内核发射，一次设备到主机的传输组成的系列。for (int i = <span style="color: #; i & <span style="color: #; ++i){cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,size, cudaMemcpyHostToDevice, stream[i]);MyKernel&&&<span style="color: #0, <span style="color: #2, <span style="color: #, stream[i]&&& (outputDevPtr + i * size, inputDevPtr + i * size, size);cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size, size, cudaMemcpyDeviceToHost, stream[i]);}　　每个流将它的hostPtr输入数组的部分拷贝到设备存储器数组inputdevPtr，调用MyKernel()内核处理inputDevPtr，然后将结果outputDevPtr传输回hostPtr同样的部分。后文描述了例子中的流如何依赖设备的计算能力重叠。必须注意为了使用重叠hostPtr必须指向分页锁定主机存储器。　　调用cudaStreamDestroy()来释放流。for (int i = <span style="color: #; i & <span style="color: #; ++i)&cudaStreamDestroy(stream[i]);&　　cudaStreamDestroy()等待指定流中所有之前的任务完成，然后释放流并将控制权返回给主机线程。
第1页：第2页：
处理 SSI 文件时出错
已有条评论
处理 SSI 文件时出错
处理 SSI 文件时出错3.1.1　CUDA运行时和驱动程序
本文所属图书&>&
《cuda专家手册：gpu编程权威指南》由英伟达公司cuda首席架构师nicholaswilt亲笔撰写，深度解析gpu的架构、系统软件、编程环境，以及cuda编程各方面的知识和各种优化技术，包含大量实用代码示例，是并行程序开发&&
CUDA运行时（经常简写为CUDART）是一个由CUDA语言集成环境使用的库。每一个版本的CUDA工具集（CUDA toolchain）都有其特定的CUDA运行时版本，应用程序会自动链接和运行时兼容的工具集，程序只有当匹配版本的CUDART在路径中可用时才能正确运行。
CUDA驱动程序是后向兼容的，支持所有同版本CUDA编写的程序和旧版本程序。它提供一个相对低层的&驱动程序API&（在cuda.h中），使开发者能够更紧凑地整理资源和为初始化定时。驱动程序的版本可由cuDriverGetVersion()函数获取。
函数返回一个十进制数，给出由驱动程序支持的CUDA版本，例如，3010代表CUDA3.1，5000代表版本5.0。
表3-1总结了CUDA每个版本中的特性。在CUDA运行时应用程序中，版本信息由cudaDeviceProp结构体中的major和minor成员提供。我们会在3.3.2小节中详细审视cudaDeviceProp。
表3-1　CUDA驱动程序每版本新特性
CUDA版本&引入的驱动程序新特性
1.1&流与事件，并发1D内存复制和内核执行
2.0&3D纹理操作
2.1&提升OpenGL互用性
2.2&可分享、映射和写结合锁定内存；等步长内存纹理操作
3.0&费米架构，多引擎复制和内核并行执行
3.1&GPUDirect
3.2&64位寻址；CUDA内核中的malloc()/free()
4.0&统一虚拟寻址；提升的线程支持；主机内存注册；GPUDirect 2.0（点对点内存复制和映射）；分层纹理
4.1&立体纹理；进程间点对点映射
4.2&开普勒架构（Kepler）
5.0&动态并行；GPUDirect RDMA
CUDA运行时需要机器中安装的驱动程序版本高于或等于由运行时支持的CUDA版本。如果驱动程序的版本低于运行时版本，CUDA应用程序会初始化失败，弹出错误cudaErrorInsufficientDriver(35)。CUDA 5.0引入了设备运行时（device runtime）概念，这是CUDA运行时的一个子集，可被CUDA内核直接调用。有关设备运行时的更详细描述在第7章给出。
您对本文章有什么意见或着疑问吗？请到您的关注和建议是我们前行的参考和动力&&
您的浏览器不支持嵌入式框架，或者当前配置为不显示嵌入式框架。}