最新帖子

页: [1] 2 3 ... 10
1
CUDA / 求问CUDA中是否支持float4的Read\Write()???
« 最后发表 作者 FDYCN 十二月 03, 2021, 05:37:37 pm »
本人常年写OpenCL,最开始是写CUDA入门的,现在回归CUDA,但没有找到相关的资料说明,cuda runtime kernel当中是否可以直接进行float4的读写和计算。举个例子,OpenCL的float4 add kernel:

__kernel void Add(__global float4* a, __global float4* b, __global float4*c){
   int tid = get_global_id(0);
   c[tid] = a[tid] + b[tid]
}//OpenCL中是合法的,或者说将入参指针变为__global float* 并在相应读写时使用vload4()\vstore4()的等价操作。

但是CUDA有类似的支持吗?比如:
__global__ void Add(float4* a, float4* b, float4*c){
   int i = threadIdx.x;
   c = a + b;
}//貌似我试了下,这样子是不合法的,那请问有没有什么和OpenCL等价的方法呢???   谢谢各位!
2
CUDA / Re: indexed constant和immediate constant的区别
« 最后发表 作者 LibAndLab 十一月 29, 2021, 07:08:20 pm »
请问这两个有什么区别呢,目前我所知道的是,constant memory是off-chip memory,每个SM都有constant cache,是on-chip memory,这个图里面的 Immediate Constant Cache和 Indexed Constant Cache和上面提到的每个SM上的constant cache有什么关系呢?什么是Immediate Constant Cache?什么是 Indexed Constant Cache?
再细化下问题:
问题1:什么是Immediate Constant Cache?什么是 Indexed Constant Cache?能否举个例子,
问题2:这个图里面的 Immediate Constant Cache和 Indexed Constant Cache和上面提到的每个SM上的constant cache有什么关系呢?
问题3:核函数传输的参数存储在constant cache中嘛?
3
CUDA / indexed constant和immediate constant的区别
« 最后发表 作者 LibAndLab 十一月 29, 2021, 06:57:07 pm »
请问这两个有什么区别呢,目前我所知道的是,constant memory是off-chip memory,每个SM都有constant cache,是on-chip memory,这个图里面的 Immediate Constant Cache和 Indexed Constant Cache和上面提到的每个SM上的constant cache有什么关系呢?什么是Immediate Constant Cache?什么是 Indexed Constant Cache?
4
NV边缘计算 / Re: Jetson AGX Xavier平台的部署工具链有么有提供效率评估工具?
« 最后发表 作者 sisiy 十月 25, 2021, 03:28:45 pm »
答复:(1)效率的理论评估没有。但是TRT有个使用随机数据进行评估的“模拟时间”,可以用这个来看看;
也可以直接TRT部署后,实际的运行一下,看看到底速度如何。而不仅仅是根据理论的,网络的结构,参数量/运算量这些来“理论”推算。直接看实际的比较方便。(也就是前提是你得买一块AGX。。。。)

(2)关于指定多个网络,每个网络分配分别哪个“处理器核心”。这个没有听说过。因为考虑到实际上GPU你没法指定“核心”的。这个似乎真心做不到。同时CPU倒是可以指定执行核心,但是Xavier的CPU是没有TensorCore之类的加速的,有TensorCore的GPU又不能按照核心指定,所以似乎并不存在一种按照核心手工指定的调度方式。
5
NV边缘计算 / Jetson AGX Xavier平台的部署工具链有么有提供效率评估工具?
« 最后发表 作者 sisiy 十月 25, 2021, 03:27:16 pm »
问:我想向您请教几个关于Jetson AGX Xavier平台的问题,非常感谢! 1、Jetson AGX Xavier平台的部署工具链有么有提供效率评估工具?例如,我有一套算法网络,可以知道网络的结构、运算量、参数量,要求能在硬件平台上100ms内推理完成。我怎么衡量Jetson AGX Xavier平台的算力能否满足需求?Jetson AGX Xavier平台有提供相应的工具吗? 2、Jetson AGX Xavier平台的软件栈或者TensorRT有没有运行时调度工具? TensorRT可以调用API指定网络在哪个处理核上跑,那是否可以运行时自动调度呢(针对多个网络场景)?
6
CUDA / Re: 关于核函数中的循环和判断
« 最后发表 作者 屠戮人神 十月 25, 2021, 03:15:22 pm »
我的意思主要有两点:
(1)代码只能看到一个骨架,最多只能看出多种计算之间是独立的,而不能知道每种计算究竟是否很简单/很复杂。如果计算很复杂的话,则可以在一定范围内分类(不一定是排序,分类即可),这样每种分类可以warp集中的计算,减少warp内分支,从而规避你担忧的因为divergent branch而导致的性能下降。

(2)你可以可以不分类整理到shared中的数组中。可以直接快速每次每个warp对一定范围内的数据(例如一个动态的,根据numpic,和行数之类的,所决定的范围),重复扫描多次,每次只处理一种。例如一个256个目标区域,第一次扫描需要处理的类型1的,warp能攒到~32个目标就集体计算一次。然后再扫描类型2,类似3的。这种不需要分配额外的分配的bucket的空间,但可能需要重复的读取原始数据多次。

(3)如果实际的问题和示范的代码不同,实际上只有有限很少量的几种,甚至一定范围内的,能落入分类1的就肯定不走范围2之类的,也可以考虑直接上,说不定更快。

提供更详细的代码有助于更进一步的分析。

此外,根据你的中文(而不是你的代码,代码并不能看出这点),你的问题实际上是一个行宽不定的数组,例如:
**
*
***
**
***
*
*
**
...

其中每个*代表一个处理任务,则等于横向有1-4个不定长度,纵向也是动态的这样的一个任务数组。

而这个动态的2D表格中的任务有多种类型。所以直接每个线程黑上一个*, 必然导致warp中的32个线程在多个类别中分支,降低效率。此时不如将一个2D的区域整体考虑,无论你是扫描多次,还是用shared中分类,这样使得每个warp能集体的处理1种类型的任务,从而尽量消除了warp内的分支,则可能有助于提升效率(得看每个分类的计算复杂度,上面已经说过了,别每种计算就一行代码,还不如你分类或者扫描的代价高,就不如直接计算了)。



7
CUDA / Re: 关于核函数中的循环和判断
« 最后发表 作者 屠戮人神 十月 25, 2021, 03:09:53 pm »
感谢您的回复,我是小白,大致理解了一下您的意思,您看一下对不对。您的意思是处理一行的数据时,可以直接进行分类,我这是8个分类,分成8个类别后,在核函数里面对这8个类别直接都进行计算,然后返回值就行。这样能够避免循环和多次判断的问题,应该是这样理解的吧?
这里我需要解释一下,因为这8个类别不是每一行我都需要进行处理,所以我才进行循环numpic次,然后对每次循环的nowpic进行判断来自哪个类别然后进行计算。所以问题回归到您说的直接就计算8个类别,我不是很清楚这样子是否能够更快?
此外,我还有一点优化方向就是,其实我实际工作中不可能有8个类别,最多可能也就4个类别,然后一行数据中大部分都只需要计算1次,少部分计算2次,最少的部分计算3次或者4次。面对这种情况,不知道您是否有更好的建议,谢谢!

我的意思主要有两点:
(1)代码只能看到一个骨架,最多只能看出多种计算之间是独立的,而不能知道每种计算究竟是否很简单/很复杂。如果计算很复杂的话,则可以在一定范围内分类(不一定是排序,分类即可),这样每种分类可以warp集中的计算,减少warp内分支,从而规避你担忧的因为divergent branch而导致的性能下降。

(2)你可以可以不分类整理到shared中的数组中。可以直接快速每次每个warp对一定范围内的数据(例如一个动态的,根据numpic,和行数之类的,所决定的范围),重复扫描多次,每次只处理一种。例如一个256个目标区域,第一次扫描需要处理的类型1的,warp能攒到~32个目标就集体计算一次。然后再扫描类型2,类似3的。这种不需要分配额外的分配的bucket的空间,但可能需要重复的读取原始数据多次。

(3)如果实际的问题和示范的代码不同,实际上只有有限很少量的几种,甚至一定范围内的,能落入分类1的就肯定不走范围2之类的,也可以考虑直接上,说不定更快。

提供更详细的代码有助于更进一步的分析。
8
CUDA / Re: 关于核函数中的循环和判断
« 最后发表 作者 残败墨叶 十月 22, 2021, 07:44:10 pm »
此外,如同我上面发的,没有必要1个线程逐步处理一行(这样不仅仅4x4B不能一次性读取,warp的每个线程之间还在大跨步的跳跃式读取,没有必要)。如果每个分支超级小,则这里也应当考虑。

当然,如果你直接上问题3, 多个线程集体将行上的多个4元组给分类了,每个线程只逐步处理一行的问题也自动不存在了(因为你将集体的处理每个bucket中的问题了)。

还有一个我需要解释给您的是,我每一行这8个分类排序其实都不是固定的,如果按照您的意思创建共享内存对每一行都预先分类和排序,那么我其实还需要创建一个数组专门用来预先保存好这每一行的中类别各自的排序是吗?然后我对8个类别进行计算的时候,就是去访问共享内存来知道我当前计算的类别排第几?我又仔细的读了一下您的回复,您的意思是我可以在共享内存中直接集体处理某一类的数据,是的,我的每个类型之间是无关的,至于您说的追加原子操作我尚未了解,我需要再进一步学习。关于共享内存的处理我还没能很好的学习,我再看看,再次感谢!
9
CUDA / Re: 关于核函数中的循环和判断
« 最后发表 作者 残败墨叶 十月 22, 2021, 07:24:44 pm »
此外,如同我上面发的,没有必要1个线程逐步处理一行(这样不仅仅4x4B不能一次性读取,warp的每个线程之间还在大跨步的跳跃式读取,没有必要)。如果每个分支超级小,则这里也应当考虑。

当然,如果你直接上问题3, 多个线程集体将行上的多个4元组给分类了,每个线程只逐步处理一行的问题也自动不存在了(因为你将集体的处理每个bucket中的问题了)。

感谢您的回复,我是小白,大致理解了一下您的意思,您看一下对不对。您的意思是处理一行的数据时,可以直接进行分类,我这是8个分类,分成8个类别后,在核函数里面对这8个类别直接都进行计算,然后返回值就行。这样能够避免循环和多次判断的问题,应该是这样理解的吧?
这里我需要解释一下,因为这8个类别不是每一行我都需要进行处理,所以我才进行循环numpic次,然后对每次循环的nowpic进行判断来自哪个类别然后进行计算。所以问题回归到您说的直接就计算8个类别,我不是很清楚这样子是否能够更快?
此外,我还有一点优化方向就是,其实我实际工作中不可能有8个类别,最多可能也就4个类别,然后一行数据中大部分都只需要计算1次,少部分计算2次,最少的部分计算3次或者4次。面对这种情况,不知道您是否有更好的建议,谢谢!
10
CUDA / Re: 两张3080之间如何实现P2P?
« 最后发表 作者 屠戮人神 十月 08, 2021, 02:21:42 pm »
我司之前的Sisiy妹子,曾经单独测试过这个问题,并且发过公众号,可能她会稍后回复你相应的公众号文章的地址。

根据核对楚星提供的信息,和Wikipedia上的Spec,3080并不支持NVLink桥了,30系列只有3090一款支持。

如果是这个样子的话,那么可能3080并不支持任何形式的P2P的Access和直接传输了,没有更好的办法。

如果你有更多发现,欢迎告诉我们。

页: [1] 2 3 ... 10