当我们使用xptxas检查寄存器使用情况时,我们会看到如下内容:
ptxas info : Used 63 registers, 244 bytes cmem[0], 51220 bytes cmem[2], 24 bytes cmem[14], 20 bytes cmem[16]
Run Code Online (Sandbox Code Playgroud)
我想知道目前是否有任何文件清楚地解释了cmem [x].将常数存储器分成多个存储体,总共有多少存储体,除了0,2,14,16之外的其他存储体用于什么?
作为旁注,@ njuffa(特别感谢你)之前在nvidia的论坛上解释过什么是0,2,14,16银行:
使用的常量存储器在常量程序"变量"(存储区1)中分区,加上编译器生成的常量(存储区14).
cmem [0]:内核参数
cmem [2]:用户定义的常量对象
cmem [16]:编译器生成的常量(其中一些可能对应源代码中的文字常量)
我想知道如何将 S3 对象作为数据成员包含在 S4 对象中,即使用组合而不是继承。这是我的代码片段。
library(randomForest)
set.seed(1337)
setClass("TestManager", slots = c(
hp = "numeric",
rfObj = "randomForest")
)
setGeneric("doIt", function(obj) standardGeneric("doIt"))
setMethod("doIt", "TestManager", function(obj) {
response <- rep(c(0, 1), times = 50) # a vector of length 100
predictors <- matrix(runif(200), nrow = 100, ncol = 2) # a matrix of dimension 100 x 2
# package "randomForest" has a function "randomForest"
# that returns an object of S3 class "randomForest"
obj@rfObj <- randomForest::randomForest(predictors, response) # …
Run Code Online (Sandbox Code Playgroud) Q1:编程指南 v11.6.0 指出以下代码模式在 Volta 及更高版本的 GPU 上有效:
if (tid % warpSize < 16) {
...
float swapped = __shfl_xor_sync(0xffffffff, val, 16);
...
} else {
...
float swapped = __shfl_xor_sync(0xffffffff, val, 16);
...
}
Run Code Online (Sandbox Code Playgroud)
为什么这样?
假设if
分支首先执行,当线程 0~15 命中该__shfl_xor_sync
语句时,它们变为非活动状态,线程 16~31 开始执行指令,直到命中相同的语句,其中前半部分和后半部分扭曲交换val
。我的理解正确吗?
如果是这样,编程指南还指出“如果目标线程处于非活动状态,则检索到的值是未定义的”并且“线程可能由于多种原因而处于非活动状态,包括……采用了与当前分支路径不同的分支路径”被亚空间处决了。” 这不是意味着if
和else
分支都会得到未定义的值吗?
Q2:在当前实现独立线程调度(Volta~Ampere)的GPU上,当if
执行分支时,不活动的线程是否仍在执行NOOP?也就是说,我是否仍应将扭曲执行视为同步执行?
Q3:同步(例如__shfl_sync
、 )是语句交错(分支中的语句 A 和 B 与分支中的 X 和 Y 交错)__ballot_sync
的唯一原因吗?我很好奇当前的 ITS 与subwarp interleaving有何不同。if
else