小编Kin*_*son的帖子

CUDA恒定存储库

当我们使用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]:编译器生成的常量(其中一些可能对应源代码中的文字常量)

cuda gpu-constant-memory

7
推荐指数
1
解决办法
2642
查看次数

S4 课程中是否可以有 S3 插槽?

我想知道如何将 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)

r s4

5
推荐指数
1
解决办法
80
查看次数

CUDA独立线程调度

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。我的理解正确吗?

如果是这样,编程指南还指出“如果目标线程处于非活动状态,则检索到的值是未定义的”并且“线程可能由于多种原因而处于非活动状态,包括……采用了与当前分支路径不同的分支路径”被亚空间处决了。” 这不是意味着ifelse分支都会得到未定义的值吗?

Q2:在当前实现独立线程调度(Volta~Ampere)的GPU上,当if执行分支时,不活动的线程是否仍在执行NOOP?也就是说,我是否仍应将扭曲执行视为同步执行?

Q3:同步(例如__shfl_sync、 )是语句交错(分支中的语句 A 和 B 与分支中的 X 和 Y 交错)__ballot_sync的唯一原因吗?我很好奇当前的 ITS 与subwarp interleaving有何不同。ifelse

cuda

5
推荐指数
1
解决办法
677
查看次数

标签 统计

cuda ×2

gpu-constant-memory ×1

r ×1

s4 ×1