标签: cuda

__syncthreads()在我有意使用线程丢弃的块中使用是否安全return？

文档说明__syncthreads() 必须由块中的每个线程调用,否则它将导致死锁,但实际上我从未经历过这样的行为.

示例代码:

__global__ void kernel(float* data, size_t size) {
    // Drop excess threads if user put too many in kernel call.
    // After the return, there are `size` active threads.
    if (threadIdx.x >= size) {
        return;
    }

    // ... do some work ...

    __syncthreads(); // Is this safe?

    // For the rest of the kernel, we need to drop one excess thread
    // After the return, there are `size - 1` active threads …

在刷新内存之前,我的CUDA程序在执行期间崩溃了.结果,设备内存仍然被占用.

我正在运行GTX 580,但nvidia-smi --gpu-reset不支持.

放置cudaDeviceReset()在程序的开头只会影响进程创建的当前上下文,并且不会刷新在它之前分配的内存.

我正在远程访问具有该GPU的Fedora服务器,因此物理重置非常复杂.

所以,问题是 - 在这种情况下有没有办法刷新设备内存？

在CUDA中,我们可以使用固定内存来更有效地将数据从主机复制到GPU,而不是通过主机上的malloc分配的默认内存.但是,默认固定存储器和零复制固定存储器有两种类型的固定存储器.

默认的固定内存将数据从主机复制到GPU的速度是正常传输速度的两倍,所以肯定有一个优势(前提是我们有足够的主机内存来锁定页面)

在固定存储器的不同版本中,即零复制存储器,我们不需要将数据从主机复制到GPU的DRAM.内核直接从主机内存中读取数据.

我的问题是:这些固定内存类型中的哪一种是更好的编程习惯.

我有2个非常相似的内核函数,在某种意义上代码几乎相同,但略有不同.目前我有2个选择:

• 写出2种不同的方法(但非常相似)
• 编写单个内核并在if/else语句中放置不同的代码块

if语句对我的算法性能有多大影响？
我知道没有分支,因为所有块中的所有线程都将输入if或else.
如果内核函数被多次调用,那么单个if语句会降低我的性能吗？

计算能力2.0的增加超过1.3,2.1超过2.0和3.0超过2.1？

我对如何在CUDA中使用共享和全局内存感到困惑,尤其是对于以下内容:

• 当我们使用时cudaMalloc(),我们得到一个指向共享或全局内存的指针吗？
• 全局内存是否驻留在主机或设备上？
• 是否有任何尺寸限制？
• 哪个访问速度更快？

• 将变量存储在共享内存中与通过内核传递其地址相同吗？也就是说,而不是

  __global__ void kernel() {
     __shared__ int i;
     foo(i);
  }
  

  Run Code Online (Sandbox Code Playgroud)

  为什么不等同呢

  __global__ void kernel(int *i_ptr) {
     foo(*i_ptr);
  }
  
  int main() {
     int *i_ptr;
     cudaMalloc(&i_ptr, sizeof(int));
     kernel<<<blocks,threads>>>(i_ptr);
  }
  

  Run Code Online (Sandbox Code Playgroud)

关于全局与共享内存中的特定速度问题存在很多问题,但没有一个问题包含何时在实践中使用任何一个内容的概述.

非常感谢

我一直在寻找C++的库/扩展,它将允许基于GPU的高级处理.我不是GPU编程方面的专家,我不想深入挖掘.我有一个由具有虚函数的类组成的神经网络.我需要一个基本上为我分配GPU的库 - 在很高的层次上.有一个人在一个名为GPU ++的系统上写了一篇论文,它为你做了大部分的GPU工作.我找不到任何地方的代码,只是他的论文.

有没有人知道类似的库,或者有没有人有GPU ++的代码？像CUDA这样的库太低了,无法处理我的大多数操作(至少在没有重写我的所有进程和算法的情况下 - 我不想这样做).

为什么没有atomicAdd()双打作为CUDA 4.0或更高版本的一部分明确实现？

从CUDA编程指南4.1的附录F第97页开始,已经实现了以下版本的atomicAdd.

int atomicAdd(int* address, int val);
unsigned int atomicAdd(unsigned int* address,
                       unsigned int val);
unsigned long long int atomicAdd(unsigned long long int* address,
                                 unsigned long long int val);
float atomicAdd(float* address, float val)

同样的页面继续为我的项目中刚刚开始使用的下面的双打提供了一个小型的atomicAdd实现.

__device__ double atomicAdd(double* address, double val)
{
    unsigned long long int* address_as_ull =
                             (unsigned long long int*)address;
    unsigned long long int old = *address_as_ull, assumed;
    do {
        assumed = old;
old = atomicCAS(address_as_ull, assumed,
                        __double_as_longlong(val +
                               __longlong_as_double(assumed)));
    } while (assumed …

我看到很多问题散布在互联网上,关于分支差异,以及如何避免分歧.然而,即使在阅读了几篇关于CUDA如何工作的文章之后,我似乎也无法看到在大多数情况下如何避免分支差异.在有人用伸出的爪子抓住我之前,请允许我描述我认为是"大多数情况".

在我看来,大多数分支差异实例涉及许多真正不同的代码块.例如,我们有以下场景:

if (A):
  foo(A)
else:
  bar(B)

如果我们有两个线程遇到这种分歧,线程1将首先执行,采用路径A.接下来,线程2将采用路径B.为了消除分歧,我们可能会将上面的块更改为如下所示:

foo(A)
bar(B)

假设foo(A)在线程2和bar(B)线程1上调用是安全的,可能会期望性能得到改善.但是,这是我看到它的方式:

在第一种情况下,线程1和2串行执行.调用这两个时钟周期.

在第二种情况下,线程1和2 foo(A)并行执行,然后bar(B)并行执行.这仍然看起来像两个时钟周期,区别在于在前一种情况下,如果foo(A)涉及从内存中读取,我想线程2可以在该延迟期间开始执行,这导致延迟隐藏.如果是这种情况,分支发散代码更快.

我对运行which nvcc和显示的不同CUDA版本感到非常困惑nvidia-smi.

我的ubuntu 16.04上安装了cuda9.2和cuda10.现在我将PATH设置为指向cuda9.2.所以当我跑:

 $ which nvcc
 /usr/local/cuda-9.2/bin/nvcc

但是,当我跑

$ nvidia-smi
Wed Nov 21 19:41:32 2018       
+-----------------------------------------------------------------------------+
| NVIDIA-SMI 410.72       Driver Version: 410.72       CUDA Version: 10.0     |
|-------------------------------+----------------------+----------------------+
| 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 106...  Off  | 00000000:01:00.0 Off |                  N/A |
| N/A   53C    P0    26W /  N/A |    379MiB /  6078MiB |      2% …