最新帖子

页: [1] 2 3 ... 10
1
CUDA / Re: 求问:有关架构方面的问题
« 最后发表 作者 屠戮人神 十二月 15, 2022, 02:07:24 pm »
1:书上所说:Fermi架构可以在每个SM上同时处理48个线程束,请问这个48个线程束是如何计算得到的?与什么参数有关呢?
2:在Fermi架构中,每个SM中含有两个线程束调度器和两个指令调度调度单元,这是不是意味着每个时间周期内有两个warp即64个线程从active thread变成计算线程开始运算操作?可这样不是和问题一中的同时处理48个线程束矛盾了?

(1)现在一般称为"Resident Warps"或者"Active Warps per SM"之类的了。最大数量一般是32或者48个warp(1024或者1536线程/SM)。用最大允许驻留的线程数量除以32得到。和计算能力本身有关(老黄决定的)。

(2)Fermi的细节我已经记不清了,太遥远了。但如果用现在的7.x和8.x来说的话,从每个周期的角度看,每个SM能发射来自4个warps的指令(有4个warp schedulers),和同时最多能驻留32个或者48个warps,并不矛盾。其他的周期中,别的warps可以被选中+发射指令的。
2
CUDA / Re: 关于共享内存的使用
« 最后发表 作者 屠戮人神 十二月 15, 2022, 02:03:32 pm »
我按照您说的问题进行了修改,现在程序可以正确运行,但是耗时却比使用全局内存还多,请问使用共享内存不一定会提高程序的运行速度吗?

使用共享内存不一定会总是提速的,这是一个正常现象。
不是特殊的银弹, 本身不具有灭杀狼人的魔法效果。
3
CUDA / 求问:有关架构方面的问题
« 最后发表 作者 皮皮虾and皮皮猪 十二月 11, 2022, 07:33:07 pm »
1:书上所说:Fermi架构可以在每个SM上同时处理48个线程束,请问这个48个线程束是如何计算得到的?与什么参数有关呢?
2:在Fermi架构中,每个SM中含有两个线程束调度器和两个指令调度调度单元,这是不是意味着每个时间周期内有两个warp即64个线程从active thread变成计算线程开始运算操作?可这样不是和问题一中的同时处理48个线程束矛盾了?
4
CUDA / Re: 关于共享内存的使用
« 最后发表 作者 queqd 十二月 07, 2022, 10:03:45 am »
你这样直接初始化了shared memory里的内容,然后大家集体就这样用,显然不行的。你需要添加同步的,在每次对shared memory的改动后,和下次读取前。

此外,同步你还不能放在if里面,你可能需要处理一下你的代码架构。

最后,看到你的下标较大,小心不要越界了。(如果你不是故意设计的如此,则需要检查你的下标使用的逻辑)
我按照您说的问题进行了修改,现在程序可以正确运行,但是耗时却比使用全局内存还多,请问使用共享内存不一定会提高程序的运行速度吗?
5
CUDA / Re: 将循环移至核函数内计算出错
« 最后发表 作者 屠戮人神 十二月 04, 2022, 11:53:14 am »
以第一种方式进行循环时可以得到正确的计算结果,但是以第二种方式循环时计算出现错误,请问是什么原因导致计算错误?
程序代码: [选择]

/*方式一*/
for (int i = 0; i < 100; i++)
{
kernel << < 100, 64 >> > (i,……);
}

void __global__ kernel(i, ……)
{
   ……
}



/*方式二*/

kernel << < 100, 64 >> > (……);

void __global__ kernel(……)
{
   

    for (int i = 0; i < 180; i++)
    {
        ……
    }
}


根据你的代码看,存在至少两种可能:
(1)前者每次kernel执行的时间较短,但启动了多次kernel。后者则只启动了1次kernel,但每次执行的时间较长。这种较长的kernel在Windows下,使用非TCC驱动的专业卡的情况下,容易引起执行超时而挂掉。

(2)两种代码并不等价,前者相对后者的每次循环,等于都在做一次"全局同步", 而后者做不到这点(你最多只能较为容易的做到block内部的线程同步,全局同步需要较为苛刻的条件,如果你真的需要这点,则反复启动kernel是最佳选择)。

(3)后者的i的范围比前者大了40%,可能会引起下标计算之类的范围问题(如果i参与了下标计算之类的式子的话)。
6
CUDA / Re: 关于共享内存的使用
« 最后发表 作者 屠戮人神 十二月 04, 2022, 11:46:05 am »
我在核函数中将全局内存赋值给共享内存进行使用,最后结果不正确,请问是赋值过程有问题还是使用过程有问题?
程序代码: [选择]
void __global__ kernel1(int i,float* d_M1, int* d_H, int* d_V, float* d_E1, float* d_E)
{
    const int n = blockDim.x * blockIdx.x + threadIdx.x;
    if (n < 2500)
    {
            extern __shared__ int s_H[];
            d_H[i * 2500 + n] = 1;
            s_H[threadIdx.x] = d_H[i*2500+n];
    for (int j = 0; j < 12; j++)
    {
            s_H[i * 2500 + n] *= d_V[(i * 2500 + n) * 12 + j];
    }

  for (int j = 0; j < 12; j++)
    {
             d_E[(i * 2500 + n) * 12 + j] = (-log(d_E1[i * 2500 + n]) - (-log(d_M1[(i * 2500 + n) * 12 + j]) / 2)) * s_H[i * 2500 + n];
    }

   }
}

你这样直接初始化了shared memory里的内容,然后大家集体就这样用,显然不行的。你需要添加同步的,在每次对shared memory的改动后,和下次读取前。

此外,同步你还不能放在if里面,你可能需要处理一下你的代码架构。

最后,看到你的下标较大,小心不要越界了。(如果你不是故意设计的如此,则需要检查你的下标使用的逻辑)

7
深度学习讨论 / Re: orin 的unified memory啥情况下要attach global和host吗
« 最后发表 作者 sisiy 十二月 03, 2022, 09:16:49 am »
答:任何Unified Memory V1.0的平台都需要的。主要包括Jetson系列,和Windows上。Linux上,Pascal+的卡不需要(V2.0平台)。Attach Host/Global/Stream的目的是明确一段Unified Memory的归属(CPU,GPU,或者GPU的部分流),因为在这些平台上,CPU和GPU,不能同时访问Unified Memory,必须手工指定一段时间的归属,然后对应的设备才能访问。否则会直接挂掉,而V2.0的平台可以直接访问,不需要attach过来,attach过去的过程。不过所有的Jetson平台,哪怕是计算能力6.0以后的,也依然属于1.0平台(实现的问题)。在他们上用Unified Memory,必须attach,或者等GPU/CPU的一者完全停下来以后(例如在某处设备同步后,CPU访问)另外一者才能访问,并不是因为有TRT的存在才需要。

嗯嗯。2.0的平台还有很多其他方面,不仅仅是同时访问不需要手工设定的问题,例如可以超量分配等等,方便了很多很多。Jetson用户如果想省事,就直接等GPU完全停下再CPU上即可。例如不要在多个线程中使用多流,或者哪怕使用了,全部流同步后再用。这样不需要attach反复折腾(但是可能运行效率稍微低点)。而Jetson平台上的最佳方案可能则是attach stream + 多流, 这种有最佳效率,但是要付出多谢几次attach调用的过程,各有利弊吧。
8
深度学习讨论 / orin 的unified memory啥情况下要attach global和host吗
« 最后发表 作者 sisiy 十二月 03, 2022, 09:14:28 am »
问题:有了解orin 的unified memory啥情况下要attach global和host吗。发现和trt放一个线程后就需要做这个,否则不需要
9
CUDA / 关于共享内存的使用
« 最后发表 作者 queqd 十一月 29, 2022, 11:06:50 am »
我在核函数中将全局内存赋值给共享内存进行使用,最后结果不正确,请问是赋值过程有问题还是使用过程有问题?
程序代码: [选择]
void __global__ kernel1(int i,float* d_M1, int* d_H, int* d_V, float* d_E1, float* d_E)
{
    const int n = blockDim.x * blockIdx.x + threadIdx.x;
    if (n < 2500)
    {
            extern __shared__ int s_H[];
            d_H[i * 2500 + n] = 1;
            s_H[threadIdx.x] = d_H[i*2500+n];
    for (int j = 0; j < 12; j++)
    {
            s_H[i * 2500 + n] *= d_V[(i * 2500 + n) * 12 + j];
    }

  for (int j = 0; j < 12; j++)
    {
             d_E[(i * 2500 + n) * 12 + j] = (-log(d_E1[i * 2500 + n]) - (-log(d_M1[(i * 2500 + n) * 12 + j]) / 2)) * s_H[i * 2500 + n];
    }

   }
}
10
CUDA / 将循环移至核函数内计算出错
« 最后发表 作者 queqd 十一月 18, 2022, 11:57:07 am »
以第一种方式进行循环时可以得到正确的计算结果,但是以第二种方式循环时计算出现错误,请问是什么原因导致计算错误?
页: [1] 2 3 ... 10