最新帖子

页: [1] 2 3 ... 10
1
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参与了下标计算之类的式子的话)。
2
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里面,你可能需要处理一下你的代码架构。

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

3
深度学习讨论 / 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调用的过程,各有利弊吧。
4
深度学习讨论 / orin 的unified memory啥情况下要attach global和host吗
« 最后发表 作者 sisiy 十二月 03, 2022, 09:14:28 am »
问题:有了解orin 的unified memory啥情况下要attach global和host吗。发现和trt放一个线程后就需要做这个,否则不需要
5
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];
    }

   }
}
6
CUDA / 将循环移至核函数内计算出错
« 最后发表 作者 queqd 十一月 18, 2022, 11:57:07 am »
以第一种方式进行循环时可以得到正确的计算结果,但是以第二种方式循环时计算出现错误,请问是什么原因导致计算错误?
7
非常感谢您的回复,确实是导致了竞态,从而导致其他函数出现数组索引越界。
8
您好,我按您说的检查了,但是发现您说的这三种情况都没有发生。我用cuda-memcheck检查了内存。总是显示是这个函数上面核函数非法访问内存的问题,请问上面的核函数出问题了,是否会造成这个函数中输出错误的这种情况呢?

你给的代码里面只有p和pos这几个坐标,用来访存。如今cuda-memcheck都报错了,还能说上面没问题(p或者pos再每个线程里面不重叠的对应数组里的具体位置?)

不妨先修正你的访存出错再说。然后再重新考虑上面三种情况(是否每个线程都拿到自己的有效位置)。
9
CUDA / Re: Invalid __global__ write of size 4
« 最后发表 作者 屠戮人神 十一月 18, 2022, 10:24:55 am »
不好意思,再次打扰了,如果使用makefile 编译的话我这样source=main.cu

RUN:=./sphgpu

CC=/data/apps/cuda/11.4/bin/nvcc

$(RUN):$(source) sph.o system.o SPHGPU.o ConfigDomain.o
   $(CC) -lineinfo $(source) sph.o system.o SPHGPU.o ConfigDomain.o -o $(RUN)
   
sph.o:sph.cu
   $(CC) -lineinfo --device-c sph.cu -o sph.o
   
system.o:system.cu
   $(CC) -lineinfo --device-c system.cu -o system.o
   
SPHGPU.o:SPHGPU.cu
   $(CC) -lineinfo --device-c SPHGPU.cu -o SPHGPU.o
   
ConfigDomain.o:ConfigDomain.cu
   $(CC) -lineinfo --device-c ConfigDomain.cu -o ConfigDomain.o
   
.PHONY:clean
clean:
   -rm -rf $(RUN)
   -rm -rf *.o
或者将 -lineinfo改成 -G之后使用cuda-memcheck还是没有出现错误的行号,麻烦您指正一下错误,我尝试在网上搜相关资料没有搜到,麻烦您了

可以考虑用cuda-gdb来确定行号,有两种方案(精确和非精确的):
(1)直接用cuda-gdb --args 你的原本程序的命令行的方式,来运行你的程序。
然后在cuda-gdb的提示符下,输入:
set cuda memcheck on
然后输入r回车继续运行你的程序。

这种情况下,运行速度会变得较慢(有的时候会非常慢),但如果有kernel访存错误,会自动停留在这一行,并显示这一行的源代码(需要-g -G)。

(2)如果方法(1)慢到了无法忍受(例如经过了很久都没有停住),可以先全速运行,然后等待kernel出错的时候,暂停住进程(而不是返回出错代码或者挂掉):
export CUDA_DEVICE_WAITS_ON_EXCEPTION=1
你原本的程序命令行 回车
(耐心等待程序挂掉)
然后当屏幕上出现类似你的kernel访存错误,请使用调试器来调试进程XXXX(一个进程编号)的时候,继续输入:
sudo cuda-gdb --pid=XXXX
然后回车几次后,会停留在出错的源代码行附近(注意不是精确的)。
如果你看到的不是源代码,而是某处含有intrinsic字样的汇编代码之类的,继续输入bt回溯,一般再上一个就是出错的你的代码的位置(的附近)。

注意第二种方法不是精确的(因为流水线之类的原因),实际位置往往会偏后一点点。不过这种方法几乎万能,稍微往上看看就能找到。
10
OpenCL / float 问题
« 最后发表 作者 lukaku 十一月 17, 2022, 05:10:45 pm »
请教下,
read_imagef 读出来的都是归一化的,[0.0, 1.0]之间的图像数据,
经过计算后,会有 1.001962, 这个其实就是1.0 ,0.00192 这个实际是0.0,  也有-0.00192,这个也是0.0
因为最小像素值1
>>> 1/255
0.00392156862745098
小于0.00392的可以忽略为0.


 如果1.001962 , 0.001962  这种值直接wrtie_imagef会有问题, 这种精度问题怎么处理,一个一个判断有点费时间,opencl write_imagef 有类似处理的方法吗?

页: [1] 2 3 ... 10