列出帖子

该操作将允许你查看该会员所有的帖子,注意你只能看到你有权限看到的板块的帖子。


显示所有帖子 - 屠戮人神

页: [1] 2 3 ... 27
1
CUDA / Re: Volta架构中规约计算的同步问题
« 于: 昨天 03:18:33 pm »
引用
for (int offset = blockDim.x / 2; offset > 0; offset /= 2)
{
    if (tid < offset)
    {
        shmem[tid] += shmem[tid + offset]; // 有问题吗?为什么?
    }
    __syncthreads();
}

同时让我们来看看你的第三个代码,这代码其实和之前的并不同,这里多了这么一行:
if (tid < offset)   <---- 注意这里的约束条件
{[tid] = [tid] + [tid + offset];}
__syncthreads();

此代码在多了这个约束条件, 在此条件下,对于同1个block内部的线程来说(假设block是1D的),
不可能在该约束条件下,存在任意2个线程,使得读取的[tid]或者[tid + offset]和另外一个线程的写入的[tid']元素相同。

这个容易证明:
如果相同,
要么(1)下标tid = tid' (tid来自于threadIdx.x, 刚才强调了是1D的,因此不可能),这是不可能的。
要么(2)下标tid + offset = tid', 即存在2个线程,他们的下标差为offset。但是刚才约束了所有的下表范围均为[0, offset), 在此范围内是找不到任意两个坐标,能坐标差为offset的(显然!)

结合(1)(2),这种是安全的。

所以你看,我们分析的时候只需要线程间的读写。而无视了同一个线程对同一个下标位置的读取,然后写入(上面说了,同一个线程这样做是安全的)。

你以后可以用这种方式分析任意类似代码。

2
CUDA / Re: Volta架构中规约计算的同步问题
« 于: 昨天 03:11:13 pm »
接上:

你的原始理解,同一个线程内部的:
引用
我不太懂是什么意思。在第二段代码中的第二行,展开的话就是
shmem[tid] = shmem[tid] + shmem[tid+16]; __syncwarp();
等号右边要读取shmem[tid],左边要写入shmem[tid]。意思是这个读取和写入的次序不能保证?

问题并不在于同一个线程内部的读取,然后写入(这个是安全的)。而在于多个线程间的读写次序。因此你之前的解释看似正确,实际上不对。这点额外强调一下。

3
CUDA / Re: Volta架构中规约计算的同步问题
« 于: 昨天 03:08:04 pm »
brucefan:

关于是否一个多个线程参与的,在某区域内的有读有写的过程是安全的,则要看:
(1)编译器能否正确的产生代码(软件读写barrier),和,
(2)硬件在执行的时候,是否能保证一定的次序。
这两点缺一不可。

我们都知道从计算能力7.0+开始,warp内不在严格的locked-step的执行了,因此需要额外引入的__syncwarp()。那么在插入了__syncwarp()的代码中,例如你给出的例子:

引用
int tid = threadIdx.x;
shmem[tid] += shmem[tid+16]; __syncwarp();
shmem[tid] += shmem[tid+8];  __syncwarp();
shmem[tid] += shmem[tid+4];  __syncwarp();
shmem[tid] += shmem[tid+2];  __syncwarp();
shmem[tid] += shmem[tid+1];  __syncwarp();

也就是, 这种常见处理1个warp最后的32个数据的规约求和的代码中,我们从warp中具体的几个线程来看来看,会出现:
线程8:[8] = [8] + [8 + 4];
线程4:[4] = [4] + [4 + 4];
然后一起__syncwarp();
则,在sync warp之前,线程8和线程4的硬件上的执行次序是无法确定的(上面说的第2点),很可能线程8已经写入改变[8]元素的值,而线程4才读取[8];(或者线程8没有改变[8]元素的值,线程[4]已经读取了[8]),

这样同一个线程4,读取到的[8]的值,存在多种可能,因此程序无法得到一个确定的结果的。(所谓的竞态,race condition, 和具体的硬件执行的时序)有关,因此是不安全的。

所以这也是为何你的下一个例子,将写入读取隔离开,额外的插入__syncwarp(), 却可以保证次序,从而结果正确的原因。

4
CUDA / Re: 关于核函数的流并发 || 核函数的嵌套并发问题
« 于: 八月 14, 2019, 01:05:18 pm »
楼上两位说的不错。我稍微补充一点。

楼主的启动线程规模太大,此时并行多个grid无可能了。你的理解也是正确的,如果这种是可能的话,那么最大驻留线程数(也就是某一时刻真正在卡上并行的线程数量)这个指标,将无意义。所以这是不可能的。

此外如同bruce说的,你的确考虑可以将1个线程的工作量加大的(注意这并不总是会提高性能)。

5
CUDA / Re: 关于CUDA分配内存的问题
« 于: 八月 12, 2019, 04:54:19 pm »
我不认为中文:“你发现当前可用....GB, 但是你无法一次性的分配到.....GB, 可以通过多次分配使用“有什么无法理解的。原谅我不能用俄语为你解释。

此外,为了防止你不小心犯错,也建议单独:
size_t bytes = sizeof(double) * height * height;

再看看bytes是不是你想要的值(常见错误)。

6
CUDA / Re: 关于CUDA分配内存的问题
« 于: 八月 12, 2019, 04:51:46 pm »
[名词2],还是不好意思打扰您一下,我尝试了查看可用内存,发现我5G的内存左右,所以很显然并不是内存不够的原因,然后您说的通过多次分配,我不太理解,是我要设定多个参数分配之后再整合的意思嘛?能否贴一下简单的代码?谢谢!

我不认为中文:“你发现当前可用....GB, 但是你无法一次性的分配到.....GB, 可以通过多次分配使用“有什么无法理解的。原谅我不能用俄语为你解释。

7
CUDA / Re: 关于CUDA分配内存的问题
« 于: 八月 09, 2019, 05:39:53 pm »
感谢您的回复,我是把bits和bytes搞混了嘛,我以为1double是8位,即8bits就是1字节,就是1bytes了。
我是在VS2010,X64环境下编译的,应该没有问题。
那么按照您的说法,我就算乘上8,这个也不会超出我内存啊,也就占据一半的样子。
您所说的可以多次分配,就是我可以对于同一个数组指针,分成前后几个语句多次分配的意思嘛?

就是说,如果你确定你的当前显存可用XX GB,例如5GB好像(见下一段),但是你可能无法直接一次性的分配到(连续的)5GB,但是如果你改成分配2GB,再分配3GB,却可能成功的这么一种现象。

具体查看当前可用显存,可以调用:
size_t available, total;
cudaMemGetInfo(&available, &total);
然后检查available的值。

8
CUDA / Re: 关于CUDA分配内存的问题
« 于: 八月 09, 2019, 05:11:12 pm »
各位[名词2]好,小白最近刚学习用CUDA,然后在用cudamalloc()分配内存的时候,我发现当我分配的大小大于一定值的时候,就会出现内存分配失败的情况。我的显卡是GTX1060,6G内存,我查询GPU内存应该是有6144Mbytes,即6442450944bytes,但是当我设定具体如下:
cudaError_t err;
   err = cudaMalloc((void **)&cuc, height*height*sizeof(double));
   if( err != cudaSuccess)
   {
       printf("CUDA error: %s\n", cudaGetErrorString(err));
   }
我本来是设定height=21600,即总共才466560000bytes,1个double是1bytes吧?为什么才这么点就会超出内存呢?
请各位前辈指教

一个double是8B,一个float是4B的。而不是1B。

注意尽量使用64-bit编译,以及,有的时候你不能一次性分配到所有的显存(但是你可以分配多次)。

9
CUDA / Re: 核函数里能不能用vector、queue等数据结构?
« 于: 八月 02, 2019, 02:27:18 pm »
目前想要在核函数里实现一个算法,中间用到了vector,sort等,找到了thrust以为可以解决问题,但是调试中发现不是这么回事,thrust只能在CPU代码上用,在核函数里面好像不能用。
所以想问一下,有没有这种库,能在核函数里面用vector、sort等功能。

你好,我不懂C++。但这些年来,需要各种容器、队列、排序等的用户,都是自行实现的。仅供参考。

10
通过nvprof显示,kernel只跑了20ms,两个Buff的传输也是20ms。但是多流同步150ms。所以这种多流只能用于Kernel运行时间特别长的吗。

一般的,你的kernel执行时间越长(或者异步传输之类的任务时间越长),则同步时候需要等待的就越久。

但我不清楚你说的这种情况(Kernel或者传输很早就完成了,同步在莫名其妙的后续占用了却很久不结束),具体的是指什么,但不妨发图看一下。

11
创建了2个stream,用来做异步测试实验:一个流传数据,kernel计算;另一个流传数据;做了一下cudaDeviceSynchronize同步,用nvprof 查看了一下,这个函数耗时150ms , 这个性能明显不能接受,是我用的时候没有注意什么细节 还是 本身就是这么耗时;PS:我用的显卡是Tesla V100-SXM2-32GB.

你如果之前设备上有kernel之类的能计算上10秒钟,你看看cudaDeviceSynchronize()卡上10秒你能不能接受。

人家就这样设计的,本身就用来同步的,此时不等上一段时间还能干啥?请原谅我说的直接。

12
CUDA / Re: GPU每次运行结果不一样是什么原因
« 于: 七月 28, 2019, 02:54:26 pm »
谢谢啦!问题已解决了,应该就是你说的那样,我限制了执行这几条命令的线程数,只有blockIdx.x==0的线程执行该命令就能实现了。
修改前:p2[tid] = p1[tid];
     p1[tid] = 1.0f;
修改后:if (blockIdx.y == 0)
   {
      p2[tid] = p1[tid];
     p1[tid] = 1.0f;
   }

嗯嗯。以及,考虑到这问题这么普遍的存在,你也许应当额外看一下手册的“原子操作”和相关章节。很多问题可以变通的用其他方式解决的。

13
CUDA / Re: GPU每次运行结果不一样是什么原因
« 于: 七月 26, 2019, 05:40:51 pm »
找到原因了,是我在核函数内给三个全局变量相互赋值出了问题,例如初始化是p1=1,p2=2,p3=3。然后后面相互赋值p3=p2;p2=p1;但最后输p3结果不是2而是1。为什么会这样的?而且我加一个中间变量来赋值也还是不行,例如swap=p2,p3=swap;swap=p1,p2=swap;结果还是和上面的一样

不妨发一下代码,例如考虑如下三个变量:
__device__ int a;
__device__ int b;
当同时有多个线程在进行tmp = a; a = b; b = tmp;的时候,因为多个GPU上的线程的执行在时间上可能是错开的,此时则会发生竞态,使得a,b并没有完成交换。这就是一种例子。

楼主不妨发一下代码,没啥可以保密的。

14
CUDA / Re: GPU每次运行结果不一样是什么原因
« 于: 七月 25, 2019, 01:29:20 pm »
代码是一样的但是每次运行结果都不一样,相差很大。感觉有点像是数值结果是一样的,但数值排列顺序每次运行出来都不一样。
想问问[名词2]们到底是什么原因?

要么是你的代码有问题,要么是某些GPU算法本身就是结果不问题的,于是产生了你说的“数值结果是一样的,但排列顺序不同”。

不妨考虑一下CPU上的一些不稳定算法,是不是有类似现象?

15
为什么同样的代码,在release模式运行结果是正确的,但在debug模式运行结果全是-4.31602e+08这个数

这是一种常见的现象,我大致列举一下常见的原因:
(A)Host部分代码有问题,在Debug下和Release下表现不一致,例如受到控制的条件编译,例如不同的文件打开路径。这些往往都是没有经过好好检查,和没有100%覆盖掉所有可能的错误代码/返回值检测造成的。
(B)Device部分代码有问题,这个比较少见,根据经常大部分都是(A)。你也应当考虑检查,例如Debug模式下kernel是否成功启动之类的原因(例如Debug下的计算能力配置并不被你的当前显卡支持,而Release却可以)

一般来说,详细的检查了所有位置的可能返回值或者错误代码,则一般可以避免。我建议楼主详细的看一下异同。

页: [1] 2 3 ... 27