我正在分析一个内核，该内核在 GTX480 中每个线程使用 25 个寄存器，每个块使用 3568 字节的共享内存。内核配置为启动 16x16 线程，线程缓存首选项设置为共享。
根据 GTX480 的规格，该设备每个 SM 有 32768 个寄存器，因此可以25 regs x 256 threads per block x 6 blocks per SM同时运行多个块。
但是，Compute Visual Profiler 和 Cuda Occupancy Calculator 报告每个 SM 只有 4 个块将处于活动状态。我想知道为什么只有 4 个块会处于活动状态，而不是我预期的 5 个。
我发现的原因是 CUDA 将使用的寄存器数量向上舍入为 26，在这种情况下，活动块的数量为 4。
为什么 CUDA 对寄存器的数量进行四舍五入？因为每个线程有 25 个寄存器，每个块有 256 个线程，所以每个 SM 最多可以有 5 个块，这显然是一个优势。
环境设置：
Device 0: "GeForce GTX 480"
CUDA Driver Version / Runtime Version 5.0 / 4.0
ptxas info: Compiling entry …允许 nvidia-smi 获取硬件级别详细信息的内部操作是什么？即使某些进程已经在 GPU 设备上运行，该工具也会执行，并获取进程的使用详细信息、名称和 ID 等。是否可以在用户级别开发这样的工具？NVML 有什么关系？
我无法让我的GLSL着色器在AMD和Nvidia硬件上运行.
我不是在寻找修复特定着色器的帮助,而是如何避免出现这些问题.是否可以检查着色器是否可以在AMD/Nvidia驱动程序上编译,而无需在具有相应硬件的机器上运行该应用程序并实际尝试它？
我知道,最后,测试是唯一可以肯定的方法,但在开发过程中我想至少避免明显的问题.
每个使用GLSL的人都必须遇到这些问题,为什么我找不到一个好方法来修复它们呢？
我在安装cuda后遇到了一个奇怪的情况...我真的按照nvidia网站建议的每一步:http://docs.nvidia.com/cuda/cuda-getting-started-guide-for-linux/# axzz3H0tm46yY ....一切似乎都很好,即使样本也可以......但是当我尝试运行Cuda中的"hello world"程序时,终端显示:
目前尚未安装"nvcc"程序.您可以通过输入以下命令安装它:sudo apt-get install nvidia-cuda-toolkit
这很奇怪,因为nvcc应该已经在nvidia网站建议的程序中安装....
有谁有想法吗？
我正在使用 Nvidia 驱动程序做一个小型 OpenCL 基准测试，我的内核执行 1024 次保险丝乘加并将结果存储在一个数组中：
#define FLOPS_MACRO_1(x) { (x) = (x) * 0.99f + 10.f; } // Multiply-add
#define FLOPS_MACRO_2(x) { FLOPS_MACRO_1(x) FLOPS_MACRO_1(x) }
#define FLOPS_MACRO_4(x) { FLOPS_MACRO_2(x) FLOPS_MACRO_2(x) }
#define FLOPS_MACRO_8(x) { FLOPS_MACRO_4(x) FLOPS_MACRO_4(x) }
// more recursive macros ...
#define FLOPS_MACRO_1024(x) { FLOPS_MACRO_512(x) FLOPS_MACRO_512(x) }
__kernel void ocl_Kernel_FLOPS(int iNbElts, __global float *pf)
{
for (unsigned i = get_global_id(0); i < iNbElts; i += get_global_size(0))
{
float f = (float) i;
FLOPS_MACRO_1024(f)
pf[i] = f;
}
} …执行CUDA脚本后,我的显卡保留内存有问题(即使使用cudaFree()).
在启动时,Total Used内存大约为128MB,但在脚本运行后,执行时内存不足.
NVIDIA-SMA:
+------------------------------------------------------+
| NVIDIA-SMI 340.29 Driver Version: 340.29 |
|-------------------------------+----------------------+----------------------+
| GPU Name Persistence-M| Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap| Memory-Usage | GPU-Util Compute M. |
|===============================+======================+======================|
| 0 GeForce GTX 660 Ti Off | 0000:01:00.0 N/A | N/A |
| 10% 43C P0 N/A / N/A | 2031MiB / 2047MiB | N/A Default |
+-------------------------------+----------------------+----------------------+
+-----------------------------------------------------------------------------+
| Compute processes: GPU Memory |
| GPU PID Process name Usage |
|=============================================================================| …我正在尝试编写OpenCL代码,利用管道等最新的OpenCL 2.0功能.到目前为止,我一直在研究AMD GPU,它们支持管道.但Nvidia驱动程序不支持OpenCL 2.0.那么有没有可用于Nvidia GPU的管道结构？我的目的是直接在两个内核之间传输数据,而不是通过全局内存传递数据.所以可以使用任何帮助我这样做的东西.
我正在按照此处的示例创建一个可变长度的本地内存数组。内核签名是这样的：
__kernel void foo(__global float4* ex_buffer,
int ex_int,
__local void *local_var)
然后我调用clSetKernelArg本地内存内核参数如下：
clSetKernelArg(*kern, 2, sizeof(char) * MaxSharedMem, NULL)
当MaxSharedMem从查询设置CL_DEVICE_LOCAL_MEM_SIZE。然后在内核内部，我将分配的本地内存拆分为多个数组和其他数据结构，并按照我认为合适的方式使用它们。所有这些都适用于 AMD（gpu 和 cpu）和 Intel 设备。但是，在 Nvidia 上，CL_INVALID_COMMAND_QUEUE当我将这个内核clFinish加入队列然后在队列上运行时，我会收到错误消息。
这是一个生成上述错误的简单内核（本地工作大小为 32）：
__kernel
void s_Kernel(const unsigned int N, __local void *shared_mem_block )
{
const ushort thread_id = get_local_id(0);
__local double *foo = shared_mem_block;
__local ushort *bar = (__local ushort *) &(foo[1000]);
foo[thread_id] = 0.;
bar[thread_id] = 0;
}
如果我在本地内存中静态分配相同的数组和数据结构，内核运行良好。有人可以为这种行为和/或解决方法提供解释吗？
我正在努力做到在Nvidia Reduction上看到的所有优化.我已经实现了前四个部分,但我在第22个幻灯片中遇到了第5部分.
我无法理解为什么提供的代码可以在没有任何syncthreads()的情况下工作的原因.线程可以访问输出中的相同内存位置.
此外,幻灯片表明,如果变量未设置为volatile,则代码将不起作用.在这方面如何变得不稳定？如果我不想调用内核,那么编程它的最佳方法是什么？
我也把这些代码放在这里供参考.
__device__ void warpReduce(volatile int* sdata, int tid) {
sdata[tid] += sdata[tid + 32];
sdata[tid] += sdata[tid + 16];
sdata[tid] += sdata[tid + 8];
sdata[tid] += sdata[tid + 4];
sdata[tid] += sdata[tid + 2];
sdata[tid] += sdata[tid + 1];
}
for (unsigned int s=blockDim.x/2; s>32; s>>=1) {
if (tid < s)
sdata[tid] += sdata[tid + s];
__syncthreads();
}
if (tid < 32) warpReduce(sdata, tid);
在此先感谢您的帮助.如果需要更多信息,请评论.
我正在阅读专业CUDA C编程,并在GPU架构概述部分:
CUDA采用单指令多线程(SIMT)架构来管理和执行32个被称为warp的组中的线程.warp中的所有线程同时执行相同的指令.每个线程都有自己的指令地址计数器和寄存器状态,并对自己的数据执行当前指令.每个SM将分配给它的线程块分区为32线程warp,然后调度它以便在可用的硬件资源上执行.
SIMT架构类似于SIMD(单指令,多数据)架构.SIMD和SIMT都通过向多个执行单元广播相同的指令来实现并行性.一个关键的区别是SIMD要求向量中的所有向量元素在一个统一的同步组中一起执行,而SIMT允许同一warp中的多个线程独立执行.即使warp中的所有线程在同一程序地址处一起启动,单个线程也可能具有不同的行为.SIMT使您能够为独立的标量线程编写线程级并行代码,以及为协调线程编写数据并行代码.SIMT模型包括SIMD不具备的三个关键功能:
➤每个线程都有自己的指令地址计数器.
➤每个线程都有自己的寄存器状态.
➤每个线程都可以有一个独立的执行路径.
第一段提到" All threads in a warp execute the same instruction at the same time.",而在第二段则提到" Even though all threads in a warp start together at the same program address, it is possible for individual threads to have different behavior.".这让我感到困惑,上述陈述似乎是矛盾的.任何人都可以解释一下吗？