最新帖子

页: [1] 2 3 ... 10
1
NV边缘计算 / Re: Nvidia AGX xavier 3D 点云显示问题
« 最后发表 作者 Lynn 五月 13, 2021, 04:08:25 pm »
没有人知道怎么弄嘛?
2
广告和有偿招募 / 根据matlab程序编写CUDA,有偿
« 最后发表 作者 bianhaiyi 五月 10, 2021, 01:11:06 pm »
自己写了一个cuda程序,结果运行时间比Matlab长。处理的数据量比较大,理论上应该加速效果明显,个人估计是图像配准处有问题。有兴趣的联系我,有偿服务。
3
NV边缘计算 / Nvidia AGX xavier 3D 点云显示问题
« 最后发表 作者 Lynn 五月 10, 2021, 10:49:07 am »
有项目需求在AGX上显示3D点云图像,尝试过Mayavi不成功,VTK不支持设备。请问有没有在AGX上实现3D点云显示的python库。谢谢
4
请教诸位:一个GPU项目,在PC上使用Quadro P620 算力6.1,跑出来11ms,移植到TX2上,耗时600+ms,这是什么原因呢?放到算力5.3的Nano上也是500+ms

查了一下资料, P620是512SP, 1.2-1.3GHz标称频率; 外加80GB/s的访存峰值. 这个和TX2并没有多少差距, 不应当慢上50多倍. 此外, TX2甚至连Nano的速度都跑不过, 这就比较诡异了.

感觉更像是测时错误, 或者正好卡在了嵌入式系统比较缓慢, 而台式机上比较快的地方. 楼主能更好的描述一下你的项目么? (而不是"一个GPU项目"这种简单的短语).
5
请教诸位:一个GPU项目,在PC上使用Quadro P620 算力6.1,跑出来11ms,移植到TX2上,耗时600+ms,这是什么原因呢?放到算力5.3的Nano上也是500+ms
6
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()做出任何评价. (当然, 你也没问).
7
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, 也请发出代码.

用代码说话, 论坛自然会进行基本的代码阅读的, 你全篇都是中文, 不见代码, 只会让事情恶化.
8
CUDA / 请教有关纹理内存的问题!!!
« 最后发表 作者 charmki 四月 20, 2021, 09:42:31 am »
大家好,想请教各位关于纹理拾取的问题:
我现在的情况是,我想用较少的线程数去处理一个已经被纹理绑定的二维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();//什么时候使用线程同步



}
9
CUDA / Re: 求教!关于数组按索引累加耗时的疑问
« 最后发表 作者 Weven 四月 19, 2021, 11:24:50 am »
补充一下,_【N】是在_global_里申请的,也就是每个线程都申请了 15个float的内存
10
CUDA / Re: 求教!关于数组按索引累加耗时的疑问
« 最后发表 作者 Weven 四月 19, 2021, 11:05:01 am »
(1)单独的一个__device__函数, 不放置在具体的__global__调用者的上下文中看, 是毫无意义的. 因为它自身不能单独启动, 也无法单独测时(你刚才说的profiler测时的顶多只能对它所在的上一层kernel代码进行实际运行测时). 甚至更连启动形状(多少个blocks, 每个block有多少个线程)都无法知道, 这些都影响时间. 结合这两个因素. 你报告的7次load读取, 和乘加操作就需要8ms的说法, 很可能是不正确的, 前提很可能不成立.

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

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

首先非常感谢您的回复!对于您提出的点,我都理解,现做出如下解释
①调用 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]答疑 :) :) :)
页: [1] 2 3 ... 10