- html - 出于某种原因,IE8 对我的 Sass 文件中继承的 html5 CSS 不友好?
- JMeter 在响应断言中使用 span 标签的问题
- html - 在 :hover and :active? 上具有不同效果的 CSS 动画
- html - 相对于居中的 html 内容固定的 CSS 重复背景?
我正在尝试使用 NVCC 选项获取有关我的 CUDA 内核中寄存器使用情况的一些信息
--ptxas-options=v
虽然全局函数一切正常,但我在设备函数方面遇到了一些困难,因为
ptxas信息:使用了N个寄存器
行。我尝试使用 noinline 关键字并将它们保存在另一个文件中,相对于调用全局函数,因为我认为 NVCC 正在报告全局函数的完整寄存器使用情况,包括被调用的设备寄存器在内联之后但没有任何变化。我可以获得有关设备功能的注册使用信息,只需将它们定义为全局即可。
你有什么建议吗?
谢谢!
最佳答案
据我了解,ptxas
(设备汇编器)仅输出其链接的代码的寄存器计数。独立的 __device__
函数不由汇编器链接,它们仅被编译。因此,汇编器不会发出设备函数的寄存器计数值。我不认为有解决方法。
但是,仍然可以通过使用 cuobjdump 从汇编器输出中转储 elf 数据来获取 __device__ 函数的寄存器占用空间。您可以按如下方式执行此操作:
$ cat vdot.cu
__device__ __noinline__ float vdot(float v1, float v2) {
return (v1 * v2);
}
__device__ __noinline__ float vdot(float2 v1, float2 v2) {
return (v1.x * v2.x) + (v1.y * v2.y);
}
__device__ __noinline__ float vdot(float4 v1, float4 v2) {
return (v1.x * v2.x) + (v1.y * v2.y) + (v1.z * v2.z) + (v1.w * v2.w);
}
$ nvcc -std=c++11 -arch=sm_52 -dc -Xptxas="-v" vdot.cu
ptxas info : 0 bytes gmem
ptxas info : Function properties for cudaDeviceGetAttribute
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Function properties for _Z4vdotff
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Function properties for cudaOccupancyMaxActiveBlocksPerMultiprocessor
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Function properties for _Z4vdot6float4S_
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Function properties for cudaMalloc
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Function properties for cudaGetDevice
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Function properties for _Z4vdot6float2S_
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Function properties for cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Function properties for cudaFuncGetAttributes
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
这里,我们在设备对象文件中拥有一组单独编译的三个 __device__
函数。在其上运行 cuobjdump 会发出大量输出,但在其中您将获得每个函数的寄存器计数:
$ cuobjdump -elf ./vdot.o
Fatbin elf code:
================
arch = sm_52
code version = [1,7]
producer = cuda
host = linux
compile_size = 64bit
compressed
<---Snipped--->
.text._Z4vdotff
bar = 0 reg = 6 lmem=0 smem=0
0xfec007f1 0x001fc000 0x00570003 0x5c980780
0x00470000 0x5c980780 0x00370004 0x5c680000
0xffe007ff 0x001f8000 0x0007000f 0xe3200000
0xff87000f 0xe2400fff 0x00070f00 0x50b00000
在设备函数输出的第二行dot(float, float)
中,您可以看到该函数使用 6 个寄存器。这是我所知道的检查设备功能寄存器足迹的唯一方法。
关于cuda - __device__ 函数中的 NVCC 寄存器使用报告,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/45957072/
CUDA 专家,如果我在主机代码中定义了一个新类型: struct float_3{ float x; float y; float z; }; 并且我已经向设备传输了一些这种类型的数据,我可以
我正在尝试在 __device__ 变量上应用内核函数,根据规范,该变量驻留在“全局内存中” #include #include "sys_data.h" #include "my_helper.c
在 http://docs.nvidia.com/cuda/cuda-c-programming-guide/#device-variable-qualifier它说 __device_ 限定符变量具
这个问题在这里已经有了答案: Check whether the code is running on the GPU or CPU (2 个回答) 7年前关闭。 我想写一些主机和设备兼容的函数,并根
有什么区别,在 CUDA 程序中定义设备常量的最佳方法是什么?在 C++,主机/设备程序中,如果我想将常量定义在设备常量内存中,我可以这样做 __device__ __constant__ float
我有一个内核,它在 if 语句中调用 device 函数。代码如下: __device__ void SetValues(int *ptr,int id) { if(ptr[threadIdx
这个问题在这里已经有了答案: External calls are not supported - CUDA (1 个回答) 关闭 7 年前。 我试图了解如何在单独的头文件中解耦 CUDA __de
如果 CUDA 内核调用 __device__ 函数, block 中进入或退出 __device__ 函数的所有线程之间是否存在任何隐式同步? 如果不是那么这意味着 block 中的一些线程可能在
我正在尝试使用 NVCC 选项获取有关我的 CUDA 内核中寄存器使用情况的一些信息 --ptxas-options=v 虽然全局函数一切正常,但我在设备函数方面遇到了一些困难,因为 ptxas信息:
在 CUDA 代码的实现过程中,我经常需要一些实用函数,这些函数将从设备和主机代码中调用。所以我将这些函数声明为 __host__ __device__。这是可以的,可能的设备/主机不兼容可以通过#i
我有一些工作代码...我在其中分配了一个设备变量指针,如下所示: float *d_var; cudaMalloc(&d_var, sizeof(float) ); 稍后在我的代码中,我想将此 var
我对我编写的 CUDA 代码的行为感到困惑。我正在为名为 DimmedGridGPU 的类中的 __device__ 函数编写测试。此类以 int DIM 为模板,我遇到的函数是返回距离输入值 x 最
我正在使用 __device__ 说明符在设备上声明一个全局变量 myvar。我没有在任何地方将它设置为有意义的值(没有像您通常那样在我的内核启动方法中使用 cudaMemcpyToSymbol)。
我有一个在主机和设备代码上都使用的类,以便更轻松地传递数据。这个类有一些操作数据的方法。一个简单的例子是: struct Vector { float x, y, z; __host__ __
在 CUDA 函数类型限定符 __device__ 和 __host__ 中可以一起使用,在这种情况下函数是为主机和设备编译的。这允许消除复制粘贴。但是,没有 __host__ __device__
我想从执行基数排序的内核内部调用独占扫描函数。但是独占扫描只需要一半的线程来完成它的工作。 独占扫描算法中需要几个 __syncthreads()。如果我在开始时有一个声明,例如 if(threadI
我不明白为什么我会收到错误 dynamic initialization is not supported for __device__, __constant__, __shared__ varia
我正在 NVIDIA CUDA GPU 上研究马尔可夫链蒙特卡罗 (MCMC) 算法实现。 CPU MCMC 算法使用高质量的 Mersenne twiner 随机数生成器,我想在我编写的 GPU 内
CUDA 有没有办法在设备端函数中动态分配内存? 我找不到任何这样做的例子。 从 CUDA C 编程手册: B.15 Dynamic Global Memory Allocation void* ma
我在很多地方读到过,__device__ 函数几乎总是由 CUDA 编译器内联。那么,当我将代码从内核移动到内核调用的__device__函数时,(通常)使用的寄存器数量不会增加,这样说是否正确? 例
我是一名优秀的程序员,十分优秀!