CUDA 7.0 速查手册

时间 2019-11-17

标签 cuda 7.0 速查手册繁體版

原文原文链接

Create by Jane/Santaizi 03:57:00 3/14/2016
All right reserved.html

速查手册基于 CUDA 7.0 toolkit documentation 并对原文进行了精简.前端

手册专一于CUDA的GPU计算方面,不涉及图形显示.如需完整档请查原文http://docs.nvidia.com/cuda/index.html#axzz42oaojUNj编程

3.2.4 Page-Locked Host Memory

在Host CPU程序中划出的内存区域供多GPU设备共享使用缓存

使用方法：并发

cudaHostAlloc() , cudaFreeHost() 分配,释放 page-locked host 内存
cudaHostRegister() page-locks 一个由 malloc 获得的内存块

优势:app

在 page-locked 的内存和 GPU 内存之间能够在 kernel 执行时异步拷贝
一些GPU设备能够直接映射 page-locked 的CPU内存，跳过拷贝步骤
在一些有 front-side bus(前端总线)的设备上, host 内存和 GPU内存能够以更高速度拷贝,用 write-combining 特性的话,速度将更快.

缺点:异步

Page-locked host 内存是稀缺资源,因此在分配时容易失败.
分配大量page-locked 内存将致使pageable 内存减小,影响整体性能.

3.2.4.1 Portable memory

在多GPU设备之间充当共享内存角色.是一个 Unified Virtual Address Space.async

使用方法:ide

cudaHostAlloc(), 传入 flag cudaHostAllocPortable
cudaHostRegister(), 传入 flag cudaHostRegisterPortable

3.2.4.2 Write-Combining Memory

默认 page-locked host 内存是以 cacheable 方式分配的.你能够用 Write-Combining 方式分配. Write-Combining 内存释放 host 的L1,L2缓存资源, 在通过PCI总线时提升最多40%的速度.函数

使用方法：

cudaHostAlloc(), 传入 flag cudaHostAllocWriteCombined

优势:

增长高速缓存的容量,使得CPU到GPU内存之间的内存拷贝加速

缺点:

从 Host 环境中读取 write-combining 内存很是慢,因此只适合 Host 往里写数据(而不读取)的状况.

3.2.4.3 Mapped Memory

host CPU内存和GPU内存之间的内存地址映射.
host 和GPU有对应的内存指针. 函数返回的是 host 指针, GPU内存指针需用 cudaHostGetDevicePointer() 获取,获取的GPU内存指针能够在 kernel中去使用.

使用方法:

cudaHostAlloc(), 传入 flag cudaHostAllocMapped
cudaHostRegister(), 传入 flag cudaHostRegisterMapped

优势:

不用在CPU-GPU之间拷贝内存数据
There is no need to use streams (see Concurrent Data Transfers) to overlap data transfers with kernel execution; the kernel-originated data transfers automatically overlap with kernel execution.

缺点:

内存映射破坏了数据的原子性, 应用程序必须使用 stream 或 events 来避免数据读写顺序控制和数据同步问题.

注意:在获取GPU内存指针以前必须使用 cudaSetDeviceFlags(), 传入 flag cudaDeviceMapHost.不然 cudaHostGetDevicePointer() 会致使错误. cudaHostGetDevicePointer() 错误也会在设备GPU不支持内存映射时产生.

属性查看:

使用设备属性 canMapHostMemory = 1(支持)查询设备支持状况.

一样注意: Atomic Functions 对映射内存的原子操做对host 和 GPU设备来讲也是非原子的.

3.2.5 Asynchronous Concurrent Execution

CUDA中如下操做是互相独立且并发的：

Host(CPU) 中的计算
Device(GPU) 中的计算
从 Host 到 Device 的数据传递
从 Device 到 Host 的数据传递
在单个 Device 内存中的数据传递
在多个 Device 内存之间的数据传递

除了Host(CPU)环境内部的数据传递是顺序同步的以外,一切和GPU有关的数据传递都是独立并发的(异步).

3.2.5.1 Concurrent Execution between Host and Device

Host 中的并发操做是经过异步函数库方法实现的,并在启动后直接移交控制权回 Host 主线程,且并不保证GPU设备已经计算完相应任务.这个模式相似于 event loop,任务在异步启动后排队等待被处理,而不阻塞主线程.下面几种操做对 host 来讲是异步调用的：

Kernel launch (kernel 函数的启动)
在单个GPU设备中的内存传递
Host 内存拷贝至 Device 内存 (64KB甚至更少的数据块传递也是异步的)
任何以 Async 为后缀的内存拷贝函数
Memory set function calls

能够设置环境变量 CUDA_LAUNCH_BLOCKING = 1 来禁止 kernel 函数的异步启动. 这个特性只能用来 debug (Notice: Debug Only!).
另外在使用 Visual Profiler Nsight 采集硬件计数器的时候 kernel 的启动也是同步的, 除非 concurrent kernel profiling 选项被开启. 以 Async 后缀的内存拷贝一样在 not page-locked 的 host 内存中是同步的.

3.2.5.2 Concurrent Kernel Execution

设备计算能力超过2.X均可以并发执行 kernel 函数. 在附录表13中可查. 不一样CUDA context中的kernel 不能并发. 使用大量 texture 和内存的 kernel 也不太可能与其余并发.

属性查看:

设备属性 concurrentKernels=1 查询设备支持状况(see Device Enumeration).

3.2.5.3 Overlap of Data Transfer and Kernel Execution

一些设备可并发执行 kernel函数和异步GPU内存拷贝操做. Host 内存块必须是 page-locked的. Device内存内部的多个内存拷贝(intra-device)和 kernal 函数甚至能够同时执行.

属性查看:

设备属性 asyncEngineCount > 0 查询设备支持状况(see Device Enumeration).
concurrentKernels = 1, 而且 asyncEngineCount > 0 查询多个Device内部内存拷贝和 kernal 的并发操做支持.

3.2.5.4 Concurrent Data Transfers

设备计算能力超过2.X 能够执行并发内存拷贝.Host 内存必须为 page-locked.

属性查看:

设备属性 asyncEngineCount = 2 查询设备支持状况(see Device Enumeration).

3.2.5.5 Streams

应用程序使用 streams 来管理上述全部并发操做.一个 stream 就是一串顺序命令. 不一样 streams 之间是乱序或同步执行的.

3.2.5.5.1 Creation and Destruction

使用方法：

下例中建立了2个 stream 并分配了一个 float array 的 page-locked 内存块给 hostPtr

cudaStream_t stream[2];
for (int i = 0; i < 2; ++i)
    cudaStreamCreate(&stream[i]);
float * hostPtr;
cudaMallocHost(&hostPtr, 2*size);

每一个 stream 都被指定顺序执行下述操做:

Host -> Device 的内存拷贝
kernel 启动
Device -> Host 的内存拷贝

for (int i = 0; i < 2; ++i)
{
cudaMemcpyAsync(inputDevPtr + isize, hostPtr + isize, size, cudaMemcpyHostToDevice, stream[i]);
MyKernel <<<100, 512, 0, stream[i]>>>(outputDevPtr + isize, inputDevPtr + isize, size);
cudaMemcpyAsync(hostPtr + isize, outputDevPtr + isize, size, cudaMemcpyDeviceToHost, stream[i]);
}

释放 streams 使用 cudaStreamDestroy().

for (int i = 0; i < 2; ++i)
    cudaStreamDestory(stream[i]);

cudaStreamDestory() 等待全部 stream 中的命令执行完毕后再销毁 stream 并返回控制权给 host 主线程,也就是说它是一个阻塞的强制同步函数.

3.2.5.5.2 Default Stream

kernel 启动和 host-device 之间的内存拷贝不须要设置特殊 stream 参数(默认设置为 0 ), 他们在stream中顺序执行.

使用方法:

使用 flag --default-stream per-thread 编译或者在 include cuda.h和cuda_runtime.h头以前定义宏 CUDA_API_PER_THREAD_DEFAULT_STREAM 那么一般 stream 将都是默认的 stream, 且每一个host 线程都有本身的 stream.
使用 flag --default-stream legacy 编译, 那么默认 stream 将会是特殊的,名叫 NULL stream ,且每一个 device 对每一个 host 线程来讲都有一个单独的 stream. NULL stream 由于它隐含的同步特性而比较特别.详细描述在 Implicit Synchronization之中
对那些没有设置 flag --default-stream 的编译来讲 --default-stream legacy 为默认的设置.

3.2.5.5.3 Explicit Synchroonization

下面列举了几种显式同步各个 streams 的方法. 为了不运算性能下降, 全部同步函数都应在须要时间控制和分离启动与内存拷贝(顺序控制)时使用.

使用方法:

cudaDeviceSynchronize() 暂停主线程并等待全部 host 线程中的 streams 中的全部命令都执行完毕,再把控制权还给主线程.
cudaStreamSynchronize() 接受一个 stream 为参数,等待该 stream 中全部命令执行完毕. 它被用来同步 host 中的某一个 stream,并容许其余 stream 异步处理.
cudaStreamWaitEvent() 接受一个 stream 和一个 event 为参数, 使得全部以后加入该 stream 的事件都等待相关 event 结束以后再开始执行. stream 参数能够为 0,代表任何命令在cudaStreamWaitEvent()执行以后,不管被加入哪一个 stream 之中都必须等待 event 结束才能开始执行.
cudaStreamQuery() 能够用来查询在某个 stream 中全部命令是否已经所有执行完毕.

3.2.5.5.4 Implicit Synchronization

若是碰到如下状况, 两个 stream 中的命令是不能并发执行的：

page-locked 的 Host 内存分配
device(GPU) 内存分配
device(GPU) 内存设置(赋值)
在同一个 Device 内存中不一样地址之间的内存拷贝
任何在 NULL stream 上的 CUDA命令
L1/shared 内存的设置切换

对于那些支持并发 kernel 执行的设备来讲, 任何操做都须要附加一个检查来查看 streamed kernel launch是否已经完成：

只有在CUDA context中全部stream 中全部 thread blocks 的kenel 启动以后才能执行.
只有在CUDA context中全部kernel 启动被确认完成以后才能执行

由于操做须要作一个 cudaStreamQuery()检查,因此为了提升性能应遵循下面两个习惯:

全部互相独立的操做应该放在非独立操做以前完成
任何形式的同步都应放到最后.

3.2.5.5.5 Overlapping Behavior

两个 stream 上的命令能够根据设备的支持状况进行重叠(并发)执行. 对于3.2.5.5.1 Creation and Destruction 例子

for (int i = 0; i < 2; ++i)
{
    cudaMemcpyAsync(inputDevPtr + i*size, hostPtr + i*size, size, cudaMemcpyHostToDevice, stream[i]);
    MyKernel <<<100, 512, 0, stream[i]>>>(outputDevPtr + i*size, inputDevPtr + i*size, size);
    cudaMemcpyAsync(hostPtr + i*size, outputDevPtr + i*size, size, cudaMemcpyDeviceToHost, stream[i]);
}

对于 stream[0]、 stream[1] 来讲,2次循环前一次中 stream[0]里的 cudaMemcpyAsync DeviceToHost 和后一次循环中 stream[1]里的 cudaMemcpyAsync HostToDevice 操做能够重叠(并发), 固然这要求设备支持并发数据传输(Concurrent Data Transfer). 可是就上述代码而言，即便设备支持并发Kernel执行(Concurrent Kernel and Kernel Execution),它也不太可能跳过两次内存拷贝过程使 stream[0]和stream[1]的 kernel执行并发,因此是隐式同步(Implicit Synchronization).为了充分利用 并发数据传输(Concurrent Data Transfer)和并发Kernel执行(Concurrent Kernel and Kernel Execution)这两个特性,重写代码以下

for (int i = 0; i < 2; ++i)
    cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,
    size, cudaMemcpyHostToDevice, stream[i]);
for (int i = 0; i < 2; ++i)
    MyKernel<<<100, 512, 0, stream[i]>>>
    (outputDevPtr + i * size, inputDevPtr + i * size, size);
for (int i = 0; i < 2; ++i)
    cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,
    size, cudaMemcpyDeviceToHost, stream[i]);

上述代码即便在设备不支持 kernel并发执行的状况下：
stream[0]的 kernel执行和 stream[1]的 cudaMemcpyAsync HostToDevice内存拷贝能够重叠, stream[0]的 cudaMemcpyAsync DeviceToHost内存操做和 stream[1]的kernel执行也能够重叠.

上述代码在设备支持 kernel并发及 data transfer并发的状况下:
stream[0] 和 stream[1]中 cudaMemcpyAsync HostToDevice/DeviceToHost 并发 ,kernel 执行并发.

两种方法比较之下后一种充分利用了设备的任务重叠并发特性(从一次增长到三次).即便设备不支持,也增长了一次重叠并发(从一次并发增长到两次).

3.2.5.5.6 Callbacks

CUDA-runtime 提供了在stream中的函数回调.

使用方法:

cudaStreamAddCallback() 若是参数传入 stream = 0 则表明等待全部在callback以前的 streams中指令完结以后函数回调.

下例添加 MyCallback函数回调至每一个 stream DeviceToHost内存拷贝操做以后:

void CUDART_CB MyCallback(cudaStream_t stream, cudaError_t status, void *data){
    printf("Inside callback %d\n", (size_t)data);
}
...
for (size_t i = 0; i < 2; ++i) {
    cudaMemcpyAsync(devPtrIn[i], hostPtr[i], size, cudaMemcpyHostToDevice,
    stream[i]);
    MyKernel<<<100, 512, 0, stream[i]>>>(devPtrOut[i], devPtrIn[i], size);
    cudaMemcpyAsync(hostPtr[i], devPtrOut[i], size, cudaMemcpyDeviceToHost,
    stream[i]);
    cudaStreamAddCallback(stream[i], MyCallback, (void*)i, 0);
}

cudaStreamAddCallback 函数最后一个参数为 0 ,是CUDA保留为了未来新功能的加入.

注意: 回调中绝对不能调用CUDA API(直接或间接), 这会致使自我调用的死循环.

3.2.5.5.7 Stream Priorities

设置 stream的优先级.

使用方法:

在建立 stream时使用 cudaStreamCreateWithPriority() 函数
使用 cudaDeviceGetStreamPriorityRange() 获取可取优先级范围 [ highest priority, lowest priority ]

例子:

// get the range of stream priorities for this device
int priority_high, priority_low;
cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high);
// create streams with highest and lowest available priorities
cudaStream_t st_high, st_low;
cudaStreamCreateWithPriority(&st_high, cudaStreamNonBlocking, priority_high);
cudaStreamCreateWithPriority(&st_low, cudaStreamNonBlocking, priority_low);

3.2.5.6 Events

events 提供了能够监控设备进程的方法.和回调同样,它在特定的 stream中被触发.
传入参数 stream = 0 表示等待全部 stream 中的命令完成后触发该事件.

3.2.5.6.1 Creation and Destruction

例子:

建立:
cudaEvent_t start, stop;
cudaEventCreat(&start);
cudaEventCreat(&stop);

销毁:
cudaEventDestroy(start);
cudaEventDestroy(stop);

3.2.5.6.2 Elapsed Time

下例使用 event 记录时间:

// 添加 start event 至全部 streams中
cudaEventRecord(start, 0);
for (int i = 0; i < 2; ++i) {
    cudaMemcpyAsync(inputDev + i * size, inputHost + i * size,
        size, cudaMemcpyHostToDevice, stream[i]);
    MyKernel<<<100, 512, 0, stream[i]>>>
        (outputDev + i * size, inputDev + i * size, size);
    cudaMemcpyAsync(outputHost + i * size, outputDev + i * size,
        size, cudaMemcpyDeviceToHost, stream[i]);
}
// 在全部命令添加完以后往 streams中添加 end event
cudaEventRecord(stop, 0);
// 同步等待全部 streams中命令完成后到达 stop event
cudaEventSynchronize(stop);
float elapsedTime;
// 记录 start event 至 stop event的时间消耗
cudaEventElapsedTime(&elapsedTime, start, stop);

3.2.5.7 Synchronous Calls

当同步函数被调用以后, 直达全部相关命令执行结束后才返回控制权.使用 cudaSetDeviceFlags() 决定在同步结束后 host 线程行为是 yield,block仍是spin.

3.2.6 Multi-Device System

3.2.6.1 Device Enumeration

一个 host 系统能够拥有多个设备Device. 例子中遍历设备并获取他们的属性.

int deviceCount;
cudaGetDeviceCount(&deviceCount);
int device;
for (device = 0; device < deviceCount; ++device)
{
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, device);
    printf("Device %d has compute capability %d.%d. \n",
        device, deviceProp.major, deviceProp.minor);
}

3.2.6.2 Device Selection

一个 Host线程能够在任什么时候候使用 cudaSetDevice() 来指配设备进行运算.并切换全部执行环境.分配内存,kernel launch,streams,events等,都在最近指定的设备GPU上运行. 若是没有指定则当前选择设备号 = 0.

例子:
size_t size = 1024sizeof(float);
cudaSetDevice(0); //切换到设备0
float p0;
cudaMalloc(&p0, size); //在设备0 上分配global内存
MyKernel<<<1000, 128>>>(p0); //在设备0 上执行kernel函数
cudaSetDevice(1); //切换到设备 1
float * p1;
cudaMalloc(&p1, size); //在设备1 上分配global内存
MyKernel<<<1000, 128>>>(p1); //在设备1 上执行kernel函数

在多GPU设备的条件下,耗时的任务能够指派给多个GPU进行运算.这是很好的.(SLI技术是多GPU完成单个任务,与这个不一样)

3.2.6.3 Stream and Event Behavior

注意: kernel launch在 stream与当前 device没有关联的状况下会失败.
失败例子:
cudaSetDevice(0); //切换到设备0
cudaStream_t s0;
cudaSreamCreate(&s0); //在当前设备0 中建立 stream s0
MyKernel<<<100,64,0,s0>>>(); //在当前设备0 中的 stream s0 中加入(异步) kernel launch指令

cudaSetDevice(1);                    //切换到设备1
cudaStream_t s1;
cudaSreamCreate(&s1);                //在当前设备1 中建立 stream s1
MyKernel<<<100,64,0,s1>>>();         //在当前设备1 中的 stream s1 中加入(异步) kernel launch指令

// 上述代码是正确的
// 下面这行代码会失败
MyKernal<<<100,64,0,s0>>>();  #Error //在当前设备1 中试图往设备0 中的 stream s0加入kernel launch指令

而内存拷贝指令却与当前设备选择无关:
// 下述代码是正确的
cudaSetDevice(0); //切换到设备0
cudaStream_t s0;
cudaSreamCreate(&s0); //在当前设备0 中建立 stream s0

cudaSetDevice(1);                    //切换到设备1
cudaMemcpyAsync(devMemPtr, hostMemPtr, size, cudaMemcpyHostToDevice, s0); //This is OK

cudaEventRecord() 在 stream与当前 device没有关联的状况下会失败.
cudaEventElapsedTime() 在 stream与当前 device没有关联的状况下会失败.

cudaEventSynchronize() , cudaEventQuery() ,cudaStreamWaitEvent() 与当前设备选择无关
所以 cudaStreamWaitEvent() 能够在多个GPU设备之间作同步.

每一个设备拥有本身的默认 stream (see Default Stream).因此不一样 GPU设备之间的任务执行是独立无序的,你须要本身控制设备间的同步问题.

3.2.6.4 Peer-to-Peer Memory Access

应用程序若是在 64位处理器上执行的话,计算能力超过2.0的 Tesla系列显卡能够互相引用他们的内存地址(i.e. 一个kernel可使用另外一个设备内存地址中的数据来执行运算) 这个点对点的内存获取特性可使用 cudaDeviceCanAccessPeer() = true检查支持状况.

点对点的内存获取功能必须使用函数 cudaDeviceEnablePeerAccess() 开启.每一个设备能够支持全局最多 8个点的内存连接.

下例为两个设备之间的数据传递:
cudaSetDevice(0);
float p0;
size_t size = 1024sizeof(float);
cudaMalloc(&p0,size);
MyKernel<<<1000,128>>>(p0);
cudaSetDevice(1);
cudaDeviceEnablePeerAccess(0,0); //开启对设备0 的点对点通道

// 在设备0 上launch kernel ,且该kernel使用设备0 中的地址 p0
MyKernel<<<1000,128>>>(p0);

3.2.6.5 Peer-to-Peer Memory Copy

两个设备之间的点对点内存拷贝.
例子:
cudaSetDevice(0); // Set device 0 as current
float* p0;
size_t size = 1024 * sizeof(float);
cudaMalloc(&p0, size); // Allocate memory on device 0
cudaSetDevice(1); // Set device 1 as current
float* p1;
cudaMalloc(&p1, size); // Allocate memory on device 1
cudaSetDevice(0); // Set device 0 as current
MyKernel<<<1000, 128>>>(p0); // Launch kernel on device 0
cudaSetDevice(1); // Set device 1 as current
cudaMemcpyPeer(p1, 1, p0, 0, size); // Copy p0 to p1
MyKernel<<<1000, 128>>>(p1); // Launch kernel on device 1
A copy (in the implicit NULL stream) between the memories of two different devices

部分掠过详细请查阅CUDA7.5 toolkit Documentation

3.2.7 Unified Virtual Address Space

当程序运行在 64位处理器上时, 一个64位的内存地址能够供全部2.0以上设备和host所使用. 全部使用 CUDA API分配的 host 内存和全部 device 内存都在这个虚拟地址范围内.(换句话说64位处理器提供的指针地址范围够大了),咱们称为这个虚拟地址为通用的(unified).
咱们称它为通用虚拟地址是由于它并不表明真实的内存地址,而是一个虚拟地址到真实地址的内存地址映射(真实的内存地址是malloc出来的内存地址),为了编程方便咱们须要多个设备和host统一使用同一个内存地址规范,而通用虚拟地址解决了这个问题.

使用 cudaPointerGetAttributes() 来判断是否内存地址是否使用了通用虚拟地址技术.
当从通用地址中读写值的时候 cudaMemcpy() 函数的参数cudaMemcpyKind 应设为flag cudaMemcpyDefault. 而且只要当前设备使用了通用地址,那么即便 host 的内存不是从CUDA API中分配的,一样也可使用(malloc/new).
经过 cudaHostAlloc() 函数分配的 host 内存直接就是使用通用地址的 page-locked 内存块(可供GPU直接读取Host内存),因此也无需使用cudaHostGetDevicePointer()来获取设备内存指针了.

优势:

使用cudaHostAlloc 分配的 page-locked 内存块将自动提高 cudaMemcpy 等拷贝函数的带宽和速度,别忘了以 cudaFreeHost 释放.
由于是 page-locked 因此GPU设备可直接读取内容.

缺点:

过多分配将下降应用程序可以使用内存,因此大多用来进行CPU和GPU之间的内存传递.

可使用设备属性 unifiedAddressing = 1查看设备是否使用了通用内存地址.

3.2.8 Interprocess Communication

全部由 host线程建立分配的 Device内存指针或者 event handle 均可以在程序进程中全部的线程使用,但不能跨进程.
若是想要跨进程使用指针和事件,必须使用 InterProcess Communication API.详细可查阅 reference manual. 而且该功能只在64位 Linux系统上受到支持.(部份内容略)

3.2.9 Error Checking

全部 run-time 函数均返回 error code.但对于异步并发(Async)函数来讲,返回错误是不可能的(基于一些缘由).因此必须使用一些 host run-time 函数来获得相关错误.

检查异步错误的惟一方法是使用对应同步函数. 使用 cudaDeviceSynchronize() 函数来同步设备已得到在设备上发生的异步错误.
你也可使用不一样级别的同步函数,好比cudaStreamSynchronize(), cudaStreamWaitEvent(), __syncthreads()等.
通常 run-time函数返回 cudaSuccess做为异常指示标志.

cudaPeekAtLastError() 用来获取错误
cudaGetLastError() 获取到错误后重置 last error = cudaSuccess.

kernel launch并不像其余 run-time函数那样返回错误标识,因此必须使用上述两种方法获取错误. 而且这两个函数必须紧跟 kernel launch函数,来得到 pre-launch errors. 由于全局只有一个Error，而咱们不但愿当中有任何函数引发的 Error 覆盖了它.为了保险起见，在 kernel launch以前也使用 cudaGetLastError()来获取以前的异常并重置为 cudaSuccess.
注意： cudaStreamQuery() 和 cudaEventQuery() 可能返回 cudaErrorNotReady ,它并不被认为是一种异常错误,因此不会被上述方法所捕捉到.

3.2.10 Call Stack

在计算能力超过2.0的设备上可使用 cudaDeviceGetLimit(), cudaDeviceSetLimit() 查询和设置调用栈的大小.
当栈溢出的时候, kernel call会失败并返回一个栈溢出错误.
数据采集自GeForce-GTX760:
cudaLimitStackSize: 1024 bytes cudaLimitPrintfFifoSize: 1048576 bytes cudaLimitMallocHeapSize: 8388608 bytes cudaLimitDevRuntimeSyncDepth: 8388608 cudaLimitDevRuntimePendingLaunchCount: 8388608

3.2.11 Texture and Surface Memory

CUDA支持一些具备 texturing功能(Tesla系列就没有)的GPU设备使用 texture 和 surface内存. 从texture 或者 surface内存中读取数据比从 global内存中读取有的优点在于如下几点：

texture 和 surface内存为读取二维数据所优化,因此在读取二维数据上能提供更高的带宽速度
地址计算由专门的计算单元进行,而无须放在 kernel中去处理.
打包的数据能够用一条指令操做来赋值给多个变量.相似于SIMD
8-bit 和 16-bit 的 integer input data 能够选择性的转换成 32-bit 的 floating-point value 于范围[0.0, 1.0] or [-1.0, 1.0]内.(一般这个功能在计算图片的颜色或灰度时十分受用)