CUDA学习(六十六)

实施限制:
动态并行性保证本文档中描述的所有语义,但某些硬件和软件资源依赖于实现,并限制使用设备运行时的程序的规模,性能和其他属性。
运行:
内存占用:
设备运行时系统软件为了各种管理目的而预留存储器,特别是用于在同步期间保存父网格状态的一个预留以及用于跟踪挂起的网格启动的第二预留。 配置控件可用于减少这些预留的大小以换取某些启动限制。 有关详细信息,请参阅下面的配置选项。
大多数保留内存被分配为父内核状态的后备存储,以便在子级启动时进行同步时使用。 保守地说,这个内存必须支持存储状态,以便在设备上实现最大数量的活动线程。 这意味着cudaDeviceSynchronize()可调用的每一代父代可能需要高达150MB的设备内存,具体取决于设备配置,即使未被全部使用,该设备配置也将不可用于程序使用。
嵌套和同步深度:
使用设备运行时,一个内核可能启动另一个内核,并且该内核可能启动另一个内核,依此类推。 每个下级发布被认为是一个新的嵌套级别,级别的总数是程序的嵌套深度。 同步深度被定义为程序在儿童启动时明确同步的最深层次。 通常这比程序的嵌套深度小一个,但是如果程序不需要在所有级别调用cudaDeviceSynchronize(),那么同步深度可能与嵌套深度明显不同。
整体最大嵌套深度限制为24,但实际上,实际限制将是系统对于每个新级别所需的内存量(请参阅上面的内存占用)。 任何导致内核处于比最大值更深的级别的启动都将失败。 请注意,这也可能适用于cudaMemcpyAsync(),它本身可能会生成内核启动。 详细信息请参阅配置选项。
默认情况下,为两个同步级别保留足够的存储空间。 可以通过调用cudaDeviceSetLimit()并指定cudaLimitDevRuntimeSyncDepth来控制此最大同步深度(并因此保留存储空间)。 为了保证成功执行嵌套程序,必须在从主机启动顶级内核之前配置要支持的级别数。 在大于指定最大同步深度的深度调用cudaDeviceSynchronize()将返回错误。
如果父内核从不调用cudaDeviceSynchronize(),系统会检测到它不需要为父级状态保留空间,那么允许进行优化。 在这种情况下,由于永远不会发生明确的父/子同步,因此程序所需的内存占用量将远小于保守最大值。 这样的程序可以指定较浅的最大同步深度以避免后备存储的过度分配。
等待内核启动:
当内核启动时,跟踪所有相关的配置和参数数据,直到内核完成。 这些数据存储在系统管理的启动池中。
启动池分为固定大小的池和性能较低的虚拟池。 设备运行时系统软件将首先尝试跟踪固定大小的池中的启动数据。 当固定大小的池已满时,虚拟池将用于跟踪新的启动。
配置选项:
设备运行时系统软件的资源分配是通过主机程序中的cudaDeviceSetLimit()API控制的。 限制必须在任何内核启动之前设置,并且在GPU正在运行程序时不得更改。
以下命名限制可能会被设置:
1
内存分配和生命周期:
cudaMalloc()和cudaFree()在主机和设备环境之间具有不同的语义。 当从主机调用时,cudaMalloc()从未使用的设备内存中分配一个新的区域。 从设备运行时调用时,这些函数映射到设备端malloc()和free()。 这意味着在设备环境中,总可分配内存仅限于设备malloc()堆大小,该大小可能比可用的未使用设备内存小。 另外,从主机程序调用cudaMalloc()分配的指针上的cudaFree()是错误的,反之亦然2
SM ID和Warp ID:
请注意,在PTX中,%smid和%warpid被定义为易失性值。 设备运行时可能会将线程块重新安排到不同的SM上,以便更有效地管理资源。 因此,在线程或线程块的生命周期中依赖%smid或%warpid保持不变是不安全的。
ECC错误:
没有任何ECC错误通知可用于CUDA内核中的代码。 整个启动树完成后,主机端会报告ECC错误。 执行嵌套程序期间出现的任何ECC错误都将生成异常或继续执行(取决于错误和配置)。
timg

评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值