列出帖子

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


显示所有帖子 - 屠戮人神

页: [1] 2 3 ... 37
1
CUDA / Re: cuda原子操作的疑问?
« 于: 六月 18, 2020, 05:57:54 pm »
回答: 直接用.
解释: int版本的也没有饱和操作(钳位到signed/unsigned的最大有效表示, 然后禁止加法), 你都能用, 为何unsigned long long int的不能用?

这就像我们普通的x86的add指令, 并没有额外的指定你的整数是signed还是unsigned的, 但这并不妨碍它进行有符号加法.

如果你真的很焦虑这点, 可以手工从PTX里导出一个.s64和.u64版本的原子加法, 但可能并无实际效果上的意义(因为用补码的整数加法来说, 有无符号并无意义, 即符号位直接当成普通的数值位即可, 无需特别处理).

2
CUDA / Re: cuda原子操作的疑问?
« 于: 六月 18, 2020, 05:47:11 pm »
在CUDA_C_Programming_Guide文件中有的atomicAdd类型有如下5种:
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);
但目前有个项目要实现   long long int +int 并得到long long int型结果  也就是我需要使用比 int 更大的数据类型long long int。 这里有什么解决办法?怎么实现呢?
先谢谢了

回答: 直接用.
解释: int版本的也没有饱和操作(钳位到signed/unsigned的最大有效表示, 然后禁止加法), 你都能用, 为何unsigned long long int的不能用?

3
你那么厉害,灌水的帖子也用AI学一学,别把自己搞的非人类了

选择不允许.cu文件参与编译, 将导致你的项目build失败.

而一个能调试的项目的前置条件是能正常build, 我不认为这个问题还有任何讨论的必要了.

4
求助!cu文件使用CUDA C/C++属性编译和使用Does not participate in build属性编译有何区别,为何在使用Does not participate in build属性编译调试时有些cuda的函数处的断点会失效,例如cudaMemcpyToSymbol、cudaBindTextureToArray。在使用CUDA C/C++编译就没有该问题,请指教,谢谢!

请勿灌水. 选择"不参与编译"将导致你的项目构建失败. 你后续的所有调试啥的都无从谈起.

此外, 你可能是一个AI, 你的"cudaMemcpyToSymbol的断点会失败", 这种说法看起来不像人类.

这是第一次警告.

5
CUDA / Re: 多个SMX间共享数据的问题
« 于: 六月 01, 2020, 01:37:52 pm »
请问各位[名词6],在最新的GPU架构中,是否实现了全局内存的一块数据只被加载一次分享给多个SMX的功能?(类似于常量内存的广播机制)

长期以来, 一直都有L2 cache, 被命中的数据将不会重复从global memory载入多次. 但除此之外, 并无其他的能"分享给多个SM"的功能.

6
OpenCL / Re: 同一环境下运行时间问题
« 于: 五月 15, 2020, 08:37:36 pm »
struct  timeval tsBegin,tsEnd;
    long t1Duration,t2Duration;
    gettimeofday(&tsBegin,NULL);
    err = clEnqueueWriteBuffer(queue,src1_memobj,CL_TRUE,0,contenLength,pHostBuffer,0,NULL,NULL);
    check_err(err,"data1 write");
    gettimeofday(&tsEnd,NULL);
    t1Duration = 1000000L*(tsEnd.tv_sec-tsBegin.tv_sec)+(tsEnd.tv_usec-tsBegin.tv_usec);
    gettimeofday(&tsBegin,NULL);
    err = clEnqueueWriteBuffer(queue,src2_memobj,CL_TRUE,0,contenLength,pHostBuffer,0,NULL,NULL);
    check_err(err,"data2 write");
    gettimeofday(&tsEnd,NULL);
    t2Duration = 1000000L*(tsEnd.tv_sec-tsBegin.tv_sec)+(tsEnd.tv_usec-tsBegin.tv_usec);

    printf("t1 duration: %ld,t2 duration:%ld\n",t1Duration,t2Duration);


t1 duration: 15,t2 duration:4

我无法为你解释μs级别的时间差异(即你的4μs和15μs).

在此微小的时间量级上, 无论是数据本身, 还是代码本身所使用的I-Cache, 或者是其他效果, 都可能会造成差异.

建议的解决方案:
(A)将数据量提升到常见的级别(例如ms级别的传输, 或者至少几百μs的), 或者,
(B)维持你的数据量不变, 但是改成连续传输三次, 使用第二次和第三次的访问/传输/API调用, 然后重新评估差异. 而不是第一次和第二次.

此外, OpenCL规范的实现比较灵活, 具体某个memory object的位置, 无法明确的确定(和你实现是否访问过这段数据, 如何访问的, 都可能有关系). 因此这应当作为一种考虑因素.

7
CUDA / Re: constant memory in multi GPU
« 于: 五月 15, 2020, 02:16:38 pm »
有一组数据,希望通过存储在常量内存中,但是如何在为多块GPU申请这样的一块constant memory?
比如:__constant__ Slice slice[num];
一共四块GPU,而且每块GPUnum的值是不一样的
只能这样申请吗?
程序代码: [选择]
#define SLICENB1 210
#define SLICENB2 234
__constant__ Slice d_slice1[SLICENB1];
__constant__ Slice d_slice2[SLICENB2];
__constant__ Slice d_slice3[SLICENB1];
__constant__ Slice d_slice4[SLICENB2];

但是在调用核函数的使用,在下面for循环中,调用的核函数的参数也就不一样了;
程序代码: [选择]
for(int i = 0; i < numGPUs; i++)
{
}

请教一下,这种情况怎么解决;如何为多块GPU分配constant memory

没有必要这样, 只要你从runtime api中, cudaSetDevice(0或1 2 3 4...)不同的设备, 然后把你的不同的值的内容, cudaMemcpyToSymbol(或者分步的等效调用)传递过来即可.

虽然这些看起来是"同名"的__constant__数组, 但是在不同的设备下, 内容可以不一样的.

进一步的解释:
实际上甚至连相同的GPU下都可以有同名的不同内容, 精确的说, 不同的CUDA Context会将这些隔离开. 你如果想知道更多内容, 可以搜索cuda context, 这点runtime api对用户进行了"隐藏"和"自动管理". 但是会在driver api中暴露出来.

8
怎么样理解在OpenCL 系统中 Kernel是如何运行起来的,比如接收到了怎样的命令的,对应的格式又是怎样的?

你好, OpenCL是一个纸面的规范, 你在具体的硬件平台上用的runtime, 则是具体的vendor对该规范的实现.

而在这个纸面的规范中, 并没有说明主机如何和设备通信, 控制设备上的kernel的启动,这并没有约定.

回到常见的GPU作为device的场合, NV对此严格保密. AMD适当披露了一些信息, 大值有将显卡的一些控制寄存器和显存, 映射成主机这边能访问的内存, 然后按照一定的格式写入该内存, 自动触发硬件上的kernel的启动发布任务.等等.

如果想知道详情, 参考AMD的ROCm(在GITHUB上), 里面有详细的描述. 我建议你看一下如何它能实现从user mode, 实现无OS/Driver的干涉的情况下, 通过往特定的内存地址, 写入特定的东西, 完成命令的发布(包括kernel的启动, 数据的传输, 等等).

9
CUDA / Re: 求最大值的问题
« 于: 五月 02, 2020, 11:18:20 pm »
如果还是按照我原来的办法做,那么数据量达到4096时,我将只能设置1024个线程,这时我考虑是1个线程去索引四个数据,但是我要想用规约是不是就需要将这些数据全部复制到shared memory中,但是我的shared memory不够大,不能存下来这么多数据,那么我该怎么做呢?

如果你的block大小是1024线程, 你完全可以这1024个线程, 每个先4个数中选出一个最大的, 然后在从这1024个最大的中继续选择.

没有人要求你必须先将这4096放入到shared中才可以, 你完全可以放1024个, 甚至512个, 256个...

不要这么死可以么? 脑子动一下啊!!

10
CUDA / Re: 求最大值的问题
« 于: 五月 02, 2020, 02:32:15 pm »
感谢回答,但是我还是有点疑问,您所说的是启动100个block,然后每个block只有一个线程,在这个单线程里面求4096个数的最大值吗?还是说开通多个线程使用规约+shared memory来计算呢?我之前是求1024个数值的最大值,我单个block里面开1024个线程,每个线程索引一个数,然后再分配1024个shared memory,利用规约去计算,这个是计算很快的,代码如下:

你能不能动动脑子? 你原本怎么将4096个数在1个block中求出最大值的? 显然是一堆线程而不是一个线程. 你显然依然这样做, 只不过将100组4096个数, 分配到100个block中而已, 这样你可以得到100个局部最大值!

没让你增加block数量就必须每个block缩小到1个线程, 你想哪里去了

11
CUDA / Re: 求最大值的问题
« 于: 五月 01, 2020, 08:47:22 pm »
假如我现在有100*4096个数,我要每4096个数求一个最大值,正常规约的思路是定义一个sharedMemory去求4096个数的最大值,但是这个4096是大于显卡所支持的最大sharedmemory的大小的,请问该怎么做呢?有什么思路吗

这东西有3个常见做法, 其中两种可以在单kernel内完成, 另外一种需要启动两次kernel. 大致如下:

(1)无kernel代码改动, 简单的将kernel启动两次即可. 第一次上100个block,每个里面依然是统计4096个数中的最大值. 这样你就得到了100个数, 其中最终的最大值必然是这100数之一. 然后第二次继续以1个block, 启动kernel, 从100数中继续选择出一个最大的.

优点: 无代码改动. 缺点: 你需要启动两次kernel, 很多人会绕不过弯.

(2)改写kernel, 只需要启动一次100个block的kernel, 每个block选出局部的一个最大值, 然后将本block的最大值用, atomicMax之类的原子操作处理后, 即可得到全体中的最大值.

优点: 一步到位, 容易理解. 缺点: 目前只有整数支持原子操作求最大值, 浮点数需要用其他方式变通实现原子操作最大值.

(3)改写kernel, 上100个block, 每个block结束前查看自己是否是最后结束的block, 如果是, 进行一次扫尾操作, 对最终的100个局部最大值完成求全局最大值.

优点: 手册上有参考实现, 可以直接抄. 缺点: 可能是最难理解的.

本论坛总是推荐方式(1).

12
CUDA / Re: 多GPU编程同步问题
« 于: 四月 27, 2020, 02:29:52 am »
问题描述:在多GPU编程中,需要在一个循环内部调用多块GPU,而且每块GPU内部也有一个for循环。在整个程序中间需要将每个GPU中间的结果输出,将这两个结构相加然后再进行新一轮的计算。将两个GPU的中间结果输出,必然需要使用同步计算。请看下面示例程序:
程序代码: [选择]
void function(int numGPUs, TGPUs* tgpu)
{
    '''
    for(int sub = 0; sub < 12; sub++)
    {
        '''
        for(int i = 0; i < numGPUs; i++)
        {
            '''
            kernelA<<<>>>(d_dataA[i], d_sum[i]);
            kernelB<<<>>>(d_dataB[i], d_sum[i]);

            cudaDeviceSynchronize(); //计算完kernelB()后需要将两块GPU中kernelB的结果进行加和“+”,然后再分别传给GPU0 和GPU1,进行下次循环
            cudaMemcpy(d_sum0[],d_sum[0], DeviceToHost);//
            cudaMemcpy(d_sum1[],d_sum[1], DeviceToHost);
            将传回主机的两个数组相加,然后再传回GPU0 和GPU1;
            //疑问1, 在这个循环里使用cudaDeviceSynchronize();当for循环循环到0块GPU时遇到cudaDeviceSynchronize()是不是就阻塞进程了,下面一个GPU需要等当前阻塞结束后才可以进行计算?
            //疑问2,cudaMemcpy()函数本身就是同步执行的,如果当前循环核函数计算完成后,使用cudaMemcpy()将结果传递到Host,是不是不用cudaDeviceSynchronize();
            //如果直接使用cudaMemcpy()直接将0块GPU的结构传递到1号GPU中进行计算,计算完成后再将最终结果分享给0号GPU,这样就需要判断当前GPU的编号,通过cudaSetDevice()标记将要进行计算的GPU。
        }
    }
}

主要问题:在for(int i = 0; i < numGPUs; i++){ ... }这个循环里面如果使用cudaDeviceSynchronize(); 或者cudaMemcpy();会不会阻塞当前循环,不能实现两块GPU共同进行计算。

希望各位接触过多GPU编程的CUDAer 都给些意见,谢谢啦~

CPU上的多线程, 并非只为了利用CPU的多个核心, 很多时候, 它可以简化逻辑. 这是很多时候被人们忽略的一点. 请思考1995年的时候, Win95推出抢占式多线程调度的时候, 距离未来的双核和多核时代还很早, 但那个时候, 只有单核心上的多线程, 依然具有很大的作用.

当你无法针对2个GPU, 分别进行普通的cudaMemcpy调用, 而非cudaMemcpyAsync系列函数调用的时候, 你应当考虑在CPU上拆分逻辑. 使用2个独立的host线程, 你可以依然保持使用简单的cudaMemcpy调用, 而无需考虑太多的其他异步或者同步调用而导致的CPU上的调度问题.

恰当的多个线程, 每个线程上的简单逻辑, 这样的组合, 等效于一个复杂的单线程的状态机. 特别当你无法好好的安排该状态机的时候, 简单的使用线程拆分逻辑, 是最后的方式. 不要陷入思维误区! 多host线程并非只是为了利用多核的CPU!

13
CUDA / Re: 流处理和cudaHostAlloc是哪个版本的cuda开始用
« 于: 四月 23, 2020, 12:54:33 pm »
谢谢

至少从10年前的CUDA 3.2时代, 甚至更早, 就可以用了. 所以你可以大胆放心使用,不必担心.

14
CUDA / Re: 小白求问:kernel核函数参数存储问题
« 于: 四月 20, 2020, 01:14:34 pm »
请问在kernel函数参数中<<<...>>>(paramA, paramB),参数paramA, paramB是存储在常量存储器还是寄存器中,我查看PTX码,都是显示的ld.param,也没看到reg的标志,

另外在核函数体中,局部变量是首先存储在寄存器中,当寄存器不够用的时候是存储在local memory中吗,(曾经好像在哪里看过local  memory是在L2缓存中,不知是否正确);

小白这两点有些疑惑,望大家解答下,跪谢 ;D ;D ;D ;D

(1)kernel的参数是存放在常量存储器中的。或者精确的说,存放在显存中,并被常量缓存缓冲。

(2)kernel中使用的任何普通变量、数组,均有可能存在放寄存器,或者存放在local memory中。这个不一定的。和当前的寄存器资源使用情况,以及,对这些变量或者数组的访问行为(例如使用指针访问地址而不是直接的名字,例如无法在编译时刻确定的下标,例如当前寄存器总数有限,使用比较紧张,而有一些变量已经很久没有使用过了等等)有关,无法直接回答一定在寄存器中或者一定在local memory中。

(3)对local memory的缓冲是在L2 cache中,还是在L1 cache(以及其各种变种),还是有其他的缓冲单位存在,和具体架构有关。

此外,PTX是一种虚拟的中间架构和对应的指令集,不适合用来观察或者猜测执行中的情况。如果你想看到实际的情况下,请使用cuobjdump --dump-sass, 而不是cuobjdump --dump-ptx来观察你的编译生成的可执行文件。前者是在具体某个架构的GPU上实际执行的指令集(和你编译时候指定的代码生成架构有关),而后者只是某种中间表示。

15
CUDA / Re: GPU cache contention问题
« 于: 四月 20, 2020, 01:06:13 pm »
在阅读GPU论文时一段warp间冲突的描述,不是很理解其意思
 Inter-warp contention happens when sibling warps contend for L1 cache space that they share and continually replace the data of each other in the cache.
注释:We call warps running on the same SIMT core sibling warps.

按照我的理解,SM每个时刻只有一个warp执行,多个warp间应该是lock-step方式。那么既然每个时刻只有一个warp在使用,为什么会出现上文所说的多个warp间竞争L1 cache呢?

(1)没有人说SM上每一个时刻(假设是周期好了)只有1个warp在执行中。
(2)没有人说SM上的warps间是“Locked Step”的执行的。
(3)所以你基于这两个不存在的前提,提出的问题也不存在。

SM上实际上总是有大量的驻留的warps(例如64个warps,或者32个或者其他数量),具体到某个具体的周期,可能有一个或者多个warps正在被调度执行(例如:4个)。这是实际情况。

Cache的使用不能只看具体到某个周期,是否有warp在访存之类的,它其中的数据的填充和淘汰,是warps们在一定的时间上(多个周期)的行为的累积效果。

结合这两段,则一个SM上有多个warps在执行,会竞争cache的说法无疑问的。你说呢。

页: [1] 2 3 ... 37