CUDA 7.0 速查手册

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

速查手册基于 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设备共享使用

使用方法：

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

优点:

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

缺点:

1. Page-locked host 内存是稀缺资源,所以在分配时容易失败.
2. 分配大量page-locked 内存将导致pageable 内存减少,影响总体性能.

3.2.4.1 Portable memory

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

使用方法:

1. cudaHostAlloc(), 传入 flag cudaHostAllocPortable
2. cudaHostRegister(), 传入 flag cudaHostRegisterPortable

3.2.4.2 Write-Combining Memory

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

使用方法：

1. cudaHostAlloc(), 传入 flag cudaHostAllocWriteCombined

优点:

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

缺点:

1. 从 Host 环境中读取 write-combining 内存非常慢,所以只适合 Host 往里写数据(而不读取)的情况.

3.2.4.3 Mapped Memory

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

使用方法:

1. cudaHostAlloc(), 传入 flag cudaHostAllocMapped
2. cudaHostRegister(), 传入 flag cudaHostRegisterMapped

优点:

1. 不用在CPU-GPU之间拷贝内存数据
2. 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.

缺点:

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

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

属性查看:

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

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

3.2.5 Asynchronous Concurrent Execution

CUDA中以下操作是互相独立且并发的：

1. Host(CPU) 中的计算
2. Device(GPU) 中的计算
3. 从 Host 到 Device 的数据传递
4. 从 Device 到 Host 的数据传递
5. 在单个 Device 内存中的数据传递
6. 在多个 Device 内存之间的数据传递

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

3.2.5.1 Concurrent Execution between Host and Device

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

1. Kernel launch (kernel 函数的启动)
2. 在单个GPU设备中的内存传递
3. Host 内存拷贝至 Device 内存 (64KB甚至更少的数据块传递也是异步的)
4. 任何以 Async 为后缀的内存拷贝函数
5. 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 也不太可能与其他并发.

属性查看:

1. 设备属性 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 函数甚至可以同时执行.

属性查看:

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

3.2.5.4 Concurrent Data Transfers

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

属性查看:

1. 设备属性 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 都被指定顺序执行下述操作:

1. Host -> Device 的内存拷贝

2. kernel 启动

3. 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中顺序执行.

使用方法:

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

3.2.5.5.3 Explicit Synchroonization

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

使用方法:

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

3.2.5.5.4 Implicit Synchronization

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

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

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

1. 只有在CUDA context中所有stream 中所有 thread blocks 的kenel 启动之后才能执行.
2. 只有在CUDA context中所有kernel 启动被确认完成之后才能执行

因为操作需要做一个 cudaStreamQuery()检查,所以为了提高性能应遵循下面两个习惯:

1. 所有互相独立的操作应该放在非独立操作之前完成
2. 任何形式的同步都应放到最后.

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中的函数回调.

使用方法:

1. cudaStreamAddCallback() 如果参数传入 stream = 0 则代表等待所有在callback之前的 streams中指令完结之后函数回调.

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

void CUDART_CB MyCallback(cudaStream_t stream, cudaError_t status, void *data){
	printf("Inside callback %d
", (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的优先级.

使用方法:

1. 在创建 stream时使用 cudaStreamCreateWithPriority() 函数
2. 使用 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. 
",
		device, deviceProp.major, deviceProp.minor);
}

3.2.6.2 Device Selection

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

例子:
size_t size = 1024*sizeof(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统一使用同一个内存地址规范,而通用虚拟地址解决了这个问题.

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

优点:

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

缺点:

1. 过多分配将降低应用程序可使用内存,所以大多用来进行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作为异常指示标志.

1. cudaPeekAtLastError() 用来获取错误
2. 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内存中读取有的优势在于以下几点：

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