列出帖子

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


显示所有帖子 - 屠戮人神

页: [1] 2 3 ... 40
1
前辈您好!我的第一个问题是想问我的event的使用是否会和数据传输起到冲突,因为cudamemcpy也存在着隐式同步,我担心它与cudaEventSynchronize()起了冲突。我对event的使用:创建event,指定cudaEventBlockingSync,然后record,最后cudaEventSynchronize(),这个流程应该没有问题吧QAQ,感谢您的回复!

我看不懂你的前半部分。。。你可以在一个流里下发cudaMemcpyAsync(),然后后面立刻随着一个event的record(你的后半段),然后等待event的record()在此流中完成。---这是你的意思?

2
所谓的线程busy polling是线程一旦调用设备(GPU)就会busy polling吗,亦或者是GPU触发同步语句包括不限于cudamecpy,才会busy polling?QAQ

你的问题的第一部分,python中的CPU(), 这个我不懂。在最初的时候说过了,欢迎其他使用python/pytorch的用户回答。

关于你的问题的第二部分,何时会spin,根据NV的官方说法:
“当前系统上的CPU核心数量,超过GPU卡的个数的时候,则在CPU<->GPU的时候,会进行忙等“。而没有提到cudaMemcpy()的时候的具体策略。

我建议总是启用使用同步对象的等待(立刻释放时间片)。你可以总是在设备级别或者Event级别启用它。

3
https://bbs.gpuworld.cn/index.php?topic=58836.0,如这篇文章所说,我想设定BlockingSync的设备标志,来阻塞host线程以获取更高的性能,否则CPU总是处于100%的自旋轮询。但是Pytorch没有这个选项,所以我考虑通过用ctypes调用cudaSetDeviceFlags,但是应该是失败,总之没有效果。pyqt的线程依然是busy polling的状态。这个问题困扰了我很久,有什么解决办法吗?不胜感激!盼回复。

我不懂python/pytorch, 但是听你的描述像是在不正确的时机设定了标志(cudaDeviceScheduleBlockingSync需要在最开头就设定好),或者你的时机正确,但是被pytorch内部给重新设定了,这个时候可能并没有什么好的办法。

一个可能的解决方案是,你额外使用ctypes对event对象设定cudaEventBlockingSync, 这个是独立的,而且可以在每次同步都选择不同的方式(例如对预计的短kernel选择spin, 或者普通kernel选择阻塞同步),而且这个可以在任意中途创建event并使用。

也欢迎你给出pytorch的原生解决方案(假设你修改了它的源代码),并发到论坛,方便其他的论坛的兄弟姐妹们。

4
CUDA / Re: cuda 矩阵计算思路求教!!!
« 于: 六月 25, 2021, 01:10:14 pm »
感谢回复,回帖不及时,抱歉!
我的M矩阵包括向量A都是cufftcomplex的浮点复数类型,之前尝试过缓冲M到sharedmemory 或寄存器中,总是会出现内存不够的情形(申请太多资源到共享内存或寄存器),因为每个元素是一个向量矩阵乘,所以可以调用cublas,但这应该属于动态并行了吧,因为程序后面会做移植,所以动态并行还没尝试过,我所用的显卡是普通的丽台P620,512个cuda核心,计算能力6.2,

(1)缓冲没有必要一次性全部缓冲完。你可以看一下手册自带的"矩阵乘法"的例子, 里面完全不是一次性缓冲完的(可以将你这个看成它的特例,想法是类似的)。
(2)如果你要缓冲M, 最坏64x64*16B(2个double), 这需要64KB, 一次性放入shared memory不可以,但是可以考虑分片,或者直接读取寄存器中, 这是可以足够放入的. 然后需要跨线程累加起来得到标量不过(直接用shuffle交换累加)。或者如果你适当的转置(和你的存储格式有关),再每次点乘,也可以不用考虑跨线程累加。

(3)从CPU端调用cublas, 和动态并行无关
(4)从设备端如果想调用,还有动态并行方面的顾虑,则现在不支持CUDA DP的显卡已经停止支持了(NV目前的Toolkit从3.5+开始支持,甚至下一个版本3.X将被彻底拿掉). 所以不用考虑这个原因。

5
CUDA / Re: 各位同事,想请教一个关于CUDA归约的问题
« 于: 六月 25, 2021, 01:01:37 pm »
你要考虑到数据和线程/blocks数量之间能否在特定的对应下,整除的问题。
例如最后多出来1个block用来处理后续的残余数据。例如可能会多出来一些线程/少出来一些数据的处理。

至于只能处理32K个,不存在这种极限。

以及,直接抄手册上的代码更快。网上这种经典代码能找出多种实现(你的两步样式的,第二步上单一block或者CPU扫尾;用原子操作一步到位的;不用原子操作单步的)。直接抄?

我举个例子说,最后最后一个对应的block:
很可能不能满足:  if(tid + 7 * blockDim.x < N)的条件(整除方面的原因)
但是会满足:
float a1 = p[tid] * re12[tid]; <--你假定的完全能进入的body
float a2 = p[tid + blockDim.x] * re12[tid + blockDim.x];
float a3 = p[tid + 2 * blockDim.x] * re12[tid + 2 * blockDim.x];
float a4 = p[tid + 3 * blockDim.x] * re12[tid + 3 * blockDim.x];
float a5 = p[tid + 4 * blockDim.x] * re12[tid + 4 * blockDim.x];
<---->实际可能a1 - a5都可以,但是条件是a8, 就会漏掉
float a6 = p[tid + 5 * blockDim.x] * re12[tid + 5 * blockDim.x];
float a7 = p[tid + 6 * blockDim.x] * re12[tid + 6 * blockDim.x];
float a8 = p[tid + 7 * blockDim.x] * re12[tid + 7 * blockDim.x];
(你假定以a8的索引条件判断)

6
CUDA / Re: 各位同事,想请教一个关于CUDA归约的问题
« 于: 六月 25, 2021, 12:56:19 pm »
在学习cuda的过程中遇到了一个有关于并行归约的困难,希望有朋友能解答一下,不胜感激。我遇到的问题如下:在展开归约时(展开因子为8),末尾的几个数据总是丢失。导致我的归约结果比正确值小一些。代码如下:主机端:
               cudaMemcpy(dev_p, p, N * sizeof(float), cudaMemcpyHostToDevice);
               cudaMemcpy(dev_re12, re12, N * sizeof(float), cudaMemcpyHostToDevice);
               RN << <blocksPerGrid/8, threadsPerBlock >> >(dev_p, dev_re12, dev_partial_rn);
               cudaMemcpy(partial_rn, dev_partial_rn, (blocksPerGrid/8) * sizeof(float), cudaMemcpyDeviceToHost);
               rn = 0;
               for (int i = 0; i<blocksPerGrid/8; i++)
                {
                    rn += partial_rn;
                }
其中threadsPerBlock=1024,blocksPerGrid=128, N=34329.
设备端:
__global__ void RN(float *p, float *re12, float *rn)
{
        __shared__ float cache[threadsPerBlock];
        int tid = threadIdx.x + blockIdx.x * blockDim.x * 8;
        int cacheIndex = threadIdx.x;
        float temp = 0;
        if(tid + 7 * blockDim.x < N)
        {       
                float a1 = p[tid] * re12[tid];
                float a2 = p[tid + blockDim.x] * re12[tid + blockDim.x];
                float a3 = p[tid + 2 * blockDim.x] * re12[tid + 2 * blockDim.x];
                float a4 = p[tid + 3 * blockDim.x] * re12[tid + 3 * blockDim.x];
                float a5 = p[tid + 4 * blockDim.x] * re12[tid + 4 * blockDim.x];
                float a6 = p[tid + 5 * blockDim.x] * re12[tid + 5 * blockDim.x];
                float a7 = p[tid + 6 * blockDim.x] * re12[tid + 6 * blockDim.x];
                float a8 = p[tid + 7 * blockDim.x] * re12[tid + 7 * blockDim.x];
                temp = a1 + a2 + a3 + a4 + a5 + a6 + a7 + a8;
        }
        cache[cacheIndex] = temp;
        // synchronize threads in this block
        __syncthreads();
        //unrolling warp
        if (blockDim.x >= 1024 && cacheIndex < 512)
        {
                cache[cacheIndex] += cache[cacheIndex + 512];
        }
        __syncthreads();
        if (blockDim.x >= 512 && cacheIndex < 256)
        {
                cache[cacheIndex] += cache[cacheIndex + 256];
        }
        __syncthreads();
        if (blockDim.x >= 256 && cacheIndex < 128)
        {
                cache[cacheIndex] += cache[cacheIndex + 128];
        }
        __syncthreads();
        if (blockDim.x >= 128 && cacheIndex < 64)
        {
                cache[cacheIndex] += cache[cacheIndex + 64];
        }
        __syncthreads();
        if (cacheIndex < 32)
        {
                volatile float *vcache = cache;
                vcache[cacheIndex] += vcache[cacheIndex + 32];
                vcache[cacheIndex] += vcache[cacheIndex + 16];
                vcache[cacheIndex] += vcache[cacheIndex + 8];
                vcache[cacheIndex] += vcache[cacheIndex + 4];
                vcache[cacheIndex] += vcache[cacheIndex + 2];
                vcache[cacheIndex] += vcache[cacheIndex + 1];
        }
        if (cacheIndex == 0)
                rn[blockIdx.x] = cache[0];
}

被这个问题困扰很久了,课题止步不前,请有能力的前辈不吝赐教,再次感谢~!

你要考虑到数据和线程/blocks数量之间能否在特定的对应下,整除的问题。
例如最后多出来1个block用来处理后续的残余数据。例如可能会多出来一些线程/少出来一些数据的处理。

至于只能处理32K个,不存在这种极限。

以及,直接抄手册上的代码更快。网上这种经典代码能找出多种实现(你的两步样式的,第二步上单一block或者CPU扫尾;用原子操作一步到位的;不用原子操作单步的)。直接抄?

7
CUDA / Re: cuda 矩阵计算思路求教!!!
« 于: 六月 15, 2021, 04:17:03 pm »
问题:计算一个960*960的图像矩阵,矩阵的每个元素是一个向量A(1*64)、矩阵M(64*64)和向量A的转置(64*1)这三者的乘积运算得来的,其中图像矩阵的每个元素对应一个向量A,根据其他输入参数得到,矩阵M为固定输入参数。
我目前的方法:一次分配对应960*960个线程,每个线程做 向量*矩阵*向量 这样的运算,然后得到结果,这样这个核函数要耗时8~9ms之久,非常不高效,
不知道大家有什么好的求解思路,感谢留言交流!

这大约是960x960x64x64x2 = 7.2G次浮点运算. 如果是最坏9ms的话, 大约是800GFlops的性能.

因为不确定你的数据类型(float or double), 具体用的卡的情况, 也不知道是否使用了常见的优化手段(例如缓冲M矩阵的整体或者部分到shared memory或者到寄存器). 因此暂时不好说你这个性能如何. 不过你总是可以考虑上cublas之类的库的(如果你不想手写的话).

还是欢迎继续讨论的.

8
CUDA / Re: 请教有关纹理内存的问题!!!
« 于: 四月 22, 2021, 02:07:22 pm »
大家好,想请教各位关于纹理拾取的问题:
我现在的情况是,我想用较少的线程数去处理一个已经被纹理绑定的二维cuda数组,假如总的线程数是256,那么纹理拾取的时候应该怎么设置x,y呢?就是texRef2D(x,y),x和y怎么和线程联系起来,并且最终可以拾取到对应位置的数组元素?

下面是我的代码:
__global__ void flowdirection(int* output, int width, int height) {//计算流向和的核函数


    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    int tid = x + y * blockDim.x * gridDim.x;//线程的绝对地址

    int DB[8] = { 1,2,4,8,16,32,64,128 };
    float flag[8] = { 0 };
    int count = 0;


    if (x < width && y < height) {

        if (tex2D(texRef, x, y) == 999) {
            return;
        }
        else {

            flag[0] = tex2D(texRef, x, y) - tex2D(texRef, x + 1, y);//1   
            flag[1] = (tex2D(texRef, x, y) - tex2D(texRef, x + 1, y + 1)) / sqrt(2.0);//2
            flag[2] = tex2D(texRef, x, y) - tex2D(texRef, x, y + 1);//4
            flag[3] = (tex2D(texRef, x, y) - tex2D(texRef, x - 1, y + 1)) / sqrt(2.0);//8
            flag[4] = tex2D(texRef, x, y) - tex2D(texRef, x - 1, y);//16
            flag[5] = (tex2D(texRef, x, y) - tex2D(texRef, x - 1, y - 1)) / sqrt(2.0);//32;
            flag[6] = tex2D(texRef, x, y) - tex2D(texRef, x, y - 1);//64
            flag[7] = (tex2D(texRef, x, y) - tex2D(texRef, x + 1, y - 1)) / sqrt(2.0);//128


            count = maxIndex(flag, 8);

            if (DB[count] > 0) {//可以解决中心像元为最低点、9个数据全为“无数据”的情况

                output[tid] = DB[count];

            }
        }




    }

    __syncthreads();//什么时候使用线程同步



}

你好, 关于你的几个问题:
(1)如果将线程的x/y坐标, 映射到问题的数据集上的x/y坐标(例如你这里的纹理的), 这个和具体问题有关, 而你已经给出写法了:
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
该x/y的映射是否合理, 需要你对着自己的问题看一下(例如你这里将线程分别按照x和y方向平铺开).
而并不存在一个固定的从blockIdx/threadIdx到(x,y,z)的映射, 这个是灵活的, 可以随意选择的.

(2)如何使用纹理拾取, 将固定线程所需要的位置的数据, 读取出来. 这个在手册上有相当多的章节进行描述, 可以自行看一下. 而你已经给出了你在使用tex2D()了. 注意也可以考虑归一化坐标, 这样线程规模和原始的数据集规模可以自动缩放, 不需要严格对应.(手册也提到了)

(3)何时在使用__syncthreads(), 以及它有什么用, 请自行阅读手册. 这里不进行解释. (以及, 你给出的代码中并不需要使用__syncthreads()的, 除非你的 maxIndex()中有对output指向的缓冲区的特殊使用).

(4)这个回复没有对你的maxIndex()做出任何评价. (当然, 你也没问).

9
CUDA / Re: 求教!关于数组按索引累加耗时的疑问
« 于: 四月 22, 2021, 01:48:13 pm »
首先非常感谢您的回复!对于您提出的点,我都理解,现做出如下解释
①调用 load_mad的 _global_函数简单来说只由两部分组成:1-计算load_mad函数参数t值   2-执行load_mad函数,如果只执行第1部分,_global_时间是μs级;
②启动形状 threads={32,16,1}   blocks={6,183, 1}
③整个_global_参数包括 para1=189*128*209(三维float,长189,宽128,高209)
 para2=189*2913*15(三维float,决定启动形状,15即为问题中循环次数N值来源) para3=189*23*2(三维float) 
             para4=428*189*12*15(四维float,load_mad函数参数p指向的数据)
④strided=1时再次验证,不影响耗时,但是如果取固定索引,如 _【i】+= t * p[100];时间是μs级

个人怀疑:_global_函数调用数据量太大?数据太多导致缓存放不下?改变启动形状会更好?
感谢[名词6]答疑 :) :) :)

如同你所说的:
当设定s=1的时候 , 读取7个数值. 也等同于固定的p[0], p[1], ...p[6]这7个元素, 不应当时间是8ms. 而只读取一个数p[200], 时间就变成了μs级了.

这可能并不是什么你想像的, 使用"索引下标"就一定会时间长或者怎么样的. 因为当时用固定的7次循环, 所有线程都访问p[0]到p[6]的时候, 下标是一致的.

此外, 这也不存在什么"缓存装不下"或者"抖动"之类的问题, 因为如果根据你现在给出的代码, 7个固定的float, 和1个固定的float, 大家都访问他们, 显存无所谓是7*4B=28字节, 还是只有1*4B=4字节的. (我们先不考虑_N[]的读取和写入)

此外, 未能理解__load_mad()是一个结构体函数的意思? 该句中文未能理解.

应当是一些你没有说明的原因, 例如p实际上和每个线程有关(例如某种推导得到的指针), 或者下标i, 并非你曾出的范例代码中的那样看上去像是线程之间uniform的. (很多人在节选代码的时候, 都丢失了细节).

如果不涉密, 请给出__global__代码. 我个人非常讨厌节选代码, 即用户自身给出"他认为问题就在这里的代码行", 而将其他行用"文字说明". 因为很可能(1)问题并不在于你给出的这几行里, (2)你文字说明或者节选掉的内容, 并不能反映真实情况. (你想想多少人心里想的一回事, 代码写出来实际却是另外一回事).

此外, 之前没说的_N[], 也未能理解你的中文"是一个按照线程分配的global memory的意思", 同样为了避免你不小心用成local, 也请发出代码.

用代码说话, 论坛自然会进行基本的代码阅读的, 你全篇都是中文, 不见代码, 只会让事情恶化.

10
CUDA / Re: 求教!关于数组按索引累加耗时的疑问
« 于: 四月 16, 2021, 01:48:21 pm »
代码如下(第一次发帖不会发图):
程序代码: [选择]
__device__ void load_mad(const f32* p, i32 s, f32 t)
{
      for (i32 i = 0; i < N; ++i)
     //for (i32 i = 0; i < 7; ++i)
     {
          _[i] += t * p[i * s];
     }
}
load_mad函数为一结构体函数,结构体变量只包括 _[N] 数组,整个函数也是在更新 _[] 数组的值。s的实际值约为127000,已测得即使s=1,也不影响函数耗时,排除耗时过高是由于显存跳跃访问带来。
其中  _[]  的总维度为N,N=15时,调用load_mad函数的核函数的耗时是18.05ms,手动设置循环次数为7时,核函数耗时为8.07ms,以上时间信息均使用NVIDIA visual Profiler测得。

我的疑问时,为什么一个看起来简单的循环计算赋值程序的耗时如此高?如何降低这种计算的耗时。
感谢各位答疑[名词6] :) :) :)

(1)单独的一个__device__函数, 不放置在具体的__global__调用者的上下文中看, 是毫无意义的. 因为它自身不能单独启动, 也无法单独测时(你刚才说的profiler测时的顶多只能对它所在的上一层kernel代码进行实际运行测时). 甚至更连启动形状(多少个blocks, 每个block有多少个线程)都无法知道, 这些都影响时间. 结合这两个因素. 你报告的7次load读取, 和乘加操作就需要8ms的说法, 很可能是不正确的, 前提很可能不成立.

(2)因为(1)你说的, 大跨步(大约隔离了128K个元素)的strided的访存并不影响性能快慢的说法, 也不成立. 这点这里不讨论(因为你的前提不成立, 无需对此"前提"下的疑问进行回复).

请修改问题, 补充信息. 或者放弃.


11
感谢回复,经过检查,发现之前统一申请内存,待所有计算完成后统一释放,改为计算完毕后,立即释放变量,这样莫名其妙就OK了,因为是新手,还是不明就里

这个和每次计算完成立刻释放内存, 和所有计算完成后统一释放无关.

需要指出的是, 频繁的每次立刻申请释放是典型的不良行为(因为cudaMalloc*()/Free()的代价高昂).

你的问题可能依然在其他地方, 只是你这样做被掩盖了.

12
源代码:
#include<stdio.h>
__global__ void k1(){
    int x=threadIdx.x%2;
    unsigned a =__ballot_sync(0x0000ffff,x);
    printf("id:%d a:%x\n",threadIdx.x,a);
}
int main(){
    k1<<<1,32>>>();
    printf("%s\n",cudaGetErrorString(cudaDeviceSynchronize()));
    return 0;
}

运行结果:
cuda-memcheck ./wallVoteFunction
========= CUDA-MEMCHECK
id:0 a:aaaa
id:1 a:aaaa
id:2 a:aaaa
id:3 a:aaaa
id:4 a:aaaa
id:5 a:aaaa
id:6 a:aaaa
id:7 a:aaaa
id:8 a:aaaa
id:9 a:aaaa
id:10 a:aaaa
id:11 a:aaaa
id:12 a:aaaa
id:13 a:aaaa
id:14 a:aaaa
id:15 a:aaaa
id:16 a:aaaa
id:17 a:aaaa
id:18 a:aaaa
id:19 a:aaaa
id:20 a:aaaa
id:21 a:aaaa
id:22 a:aaaa
id:23 a:aaaa
id:24 a:aaaa
id:25 a:aaaa
id:26 a:aaaa
id:27 a:aaaa
id:28 a:aaaa
id:29 a:aaaa
id:30 a:aaaa
id:31 a:aaaa
no error
========= ERROR SUMMARY: 0 errors

未参与投票的线程使用__ballot_sync()返回值会有什么问题吗?

你好, 不在mask中的线程, 请:
(1)不要参与ballot.
(2)不要使用ballot的结果.
(例如, 被某种条件给分支掉或者predicated off的线程)

否则可能会产生未定义的结果(例如你这个代码). 该未定义行为可能在部分条件下, 和部分计算能力的卡上, 产生某种特定的看似"有意义"的结果, 但这种结果很可能在, 以后别的条件下, 或者别的计算能力的卡上, 变得无意义.

13
如题,求教如下:
 
    float* d_x;
    float* d_y;
    float* d_px;
    float* d_py;
    float* d_input;
    float* d_output;

    Kernel << <BlocksPerGrid, ThreadsPerBlock>> > (d_x, d_y,d_px, d_py,d_input, d_output);

该核函数中的前四个数据传递到核函数中,其数值都发生了改变,而第五个参数d_input很正确的传递了进去,请问这是什么原因,有朋友遇到过吗?

楼主你好, 根据你的代码, 你的这5个参数, 均是指针. 只有对指针的进一步提领(de-reference), 才能得到具体的float数值. 所以你的描述, 前4个参数本身"接近于0.0", 这种描述是无意义的.

可能的原因是, 你对这些指针的使用有误(例如错误的偏移量或者地址值变换), 从而访问到错误的数据.

请重新检查情况, 如果有必要, 重新描述问题.

14
CUDA / Re: GPU程序一次循环后怎么清空显存?
« 于: 三月 15, 2021, 04:22:58 pm »
我的程序在main这有一个大循环,循环5次后提示内存不足,用nvidia-smi查看时发现每次循环的过程中显存都在累加,我想在每次循环后清空显存,有什么办法嘛?

根据你的上下文, 你需要用完资源后释放.

请搜索cudaFree()字样以学习如何释放不再需要的显存.

15
OpenCL / Re: 请教关于OpenCL Kernel launch的过程
« 于: 二月 23, 2021, 01:09:57 pm »
想请教论坛里的前辈们,OpenCL的kernel是如何在板子加载起来的,之前看到过一个贴子说AMD在这方面透露的一点”将显卡的一些控制寄存器和显存映射成主机这边可以访问的内存“,想请教前辈们这部分的内容在ROCm上有吗,能在哪里找到再详细一点的介绍?

有, AMD现在的驱动是开源的, 请直接查看ROCk(驱动)和ROCr的相关内容.

也可以搜索一下HSA相关字样, 看来一下无API的kernel启动方式(将一些GPU的寄存器mmio到用户地址空间, 然后通过对它们的读写, 完成无API调用的kernel启动. 也包括可以查看将DMA引擎映射到用户地址空间, 也无API的发起传输, 等等).

页: [1] 2 3 ... 40