列出帖子

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


显示所有帖子 - 屠戮人神

页: [1] 2 3 ... 45
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
您好,我按您说的检查了,但是发现您说的这三种情况都没有发生。我用cuda-memcheck检查了内存。总是显示是这个函数上面核函数非法访问内存的问题,请问上面的核函数出问题了,是否会造成这个函数中输出错误的这种情况呢?

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

不妨先修正你的访存出错再说。然后再重新考虑上面三种情况(是否每个线程都拿到自己的有效位置)。

4
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回溯,一般再上一个就是出错的你的代码的位置(的附近)。

注意第二种方法不是精确的(因为流水线之类的原因),实际位置往往会偏后一点点。不过这种方法几乎万能,稍微往上看看就能找到。

5
CUDA / Re: 从Device端传回Host端的时候cudaMemcpy函数不起作用
« 于: 十一月 03, 2022, 12:02:20 pm »
非常感谢您的回复!那我们先解决上面的问题!

前面的i,也是3203,与后面的p1=3202是一样的值,两部分代码中的sortpart[p1]以及sortpart均为0,我觉得这里应该没有问题。

您说的第二个问题是这个意思吗:在分配d_RepForce[](用于GPU端的排斥力变量)的空间时换成cudaMallocManaged(),以及不用写cudaMemcpy回传代码,直接Printf就可以看到在GPU计算后改变的值对吧?

如果第一点你能确保就好。

第二点是你这样的理解的(如果你用Windows,请确保CPU在使用的时候,GPU已经所有工作都完成了,没有在进行中;Linux可以暂时无视这点)。注意将不再存在d_RepForce和RepForce, 只有1个缓冲区,CPU和GPU都用。

6
CUDA / Re: 从Device端传回Host端的时候cudaMemcpy函数不起作用
« 于: 十一月 03, 2022, 11:52:14 am »
还有一个问题是,我在这个核函数刚开始写的赋值代码
 if(sortpart[p1] >= 0 && sortpart[p1] <CONSTANTS.nBound2Start)
{
Repforce[sortpart[p1]].x=0.0;
Repforce[sortpart[p1]].y=0.0;
Repforce[sortpart[p1]].z=0.0;
}
对最后的计算结果影响很大,加上他就是对的,不加就是错的,但是我明明在CPU端已经给他赋值为0了再传回到GPU端上,按说在这里是否赋值为0没有影响才对的吧

这是另外一个问题了,我们建议每次只考虑一个问题,并进行讨论。否则容易歪楼。

如果你能遇到这种情况,排除了没有正确初始化/传输的可能后,则较大的可能则是你有多个索引值p(或者p1), 能进行sortpart[p]这种查表映射后,映射到相同的位置,从而引发GPU上的竞态。我建议你先解决正事再讨论这个。这里只是简单说一说。你也可以自行排查一下(我看不到该表,不能确定或者排除此可能)。

7
CUDA / Re: 新手求教CPU与GPU数据储存
« 于: 十一月 03, 2022, 11:49:09 am »
请问一下,我在CPU的C++程序(大型程序中间的一步)中使用double型变量储存并输出1.250000-0.750000得到的值为0.500018 ,如果写一个测试的c++小程序,只计算这个减法不管是double或float都输出的是0.500000  ,但是在GPU端(大型程序中的一步)用float型变量得到的是0.500000 ,怎么结果会差这么多呢

考虑到你给出的例子中的两个数值(1.25和0.75), 正巧在二进制小数里面,都可以被精确表示(1.01'b和0.11'b),他们两个的相减,并不存在任何in-exact的情况,必定会得到精确的0.5的(这个数值也可以被精确表示)。

所以你在某个大型的程序中发现运算得到了其他值,而单纯的提取出来在CPU和GPU算都能得到精确值,则必然是你的"大型程序"中途混杂了其他运算步骤,干扰了你对结果的观测,这是我感觉可能的唯一理由。

8
CUDA / Re: 从Device端传回Host端的时候cudaMemcpy函数不起作用
« 于: 十一月 03, 2022, 11:41:32 am »
(1)

CPU端的检查代码:

程序代码: [选择]
for(int i=0;i<nFluidPar;i++)
{
    if(SortPart[i]==0)
        printf("%d, %f, %f, %f aaaaa\n ",i,RepForce[SortPart[i]].x,RepForce[SortPart[i]].y,RepForce[SortPart[i]].z);
}

GPU端的检查代码:
程序代码: [选择]
if(p1==3203)
      printf("%d, %f,%f ,%f BBBB\n ",sortpart[p1],Repforce[sortpart[p1]].x,Repforce[sortpart[p1]].y,Repforce[sortpart[p1]].z);

这两个似乎不太等价,可能会对你的printf的结果的查看造成干扰,前面的%d输出的是i的索引,后面的%d对应的则是sortpart[索引]变换后的。这里会不会正好对你造成误会?


(2)此外,将你的cudaMalloc改成cudaMallocManaged(), 然后取消掉最后一个cudaMemcpy回传,看看结果是否正确。Unified Memory的GPU写入结果,CPU可以直接看到(有一定的限制条件,这里先不管),不需要回传,这样可以排除或者确认是否是cudaMemcpy()本身的问题。



9
深度学习讨论 / Re: 线程束分化
« 于: 十一月 03, 2022, 09:58:00 am »
想问下支持独立线程调度的卡指的是SIMT工作模式的卡吗?也就是说在新卡上可能存在这种情况:一个wrap里所有线程执行的指令可能并不一样?

是从计算能力7.0开始引入的东西,不是SIMT。你可以参考一下手册的这里:

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#independent-thread-scheduling-7-x

10
CUDA / Re: 从Device端传回Host端的时候cudaMemcpy函数不起作用
« 于: 十一月 03, 2022, 09:54:18 am »
您好,我在这个核函数里printf了确实是计算了,以及在调用这个核函数之后用了cudaGetLastError();仍然显示no error,您看这是怎么回事呢
void CalFluidRepForce(int n, double3* TotalCoorg, int3* CCg, int* nCubeg, int* SortPartg, float* OuterReigong, int2* BeginEndCubeg, float3* RepForceg)
   {
      unsigned gridesize = unsigned(n + SPHSIZE - 1) / SPHSIZE;
      KerCalFluidRepForce << <gridesize, SPHSIZE >> >(n, TotalCoorg, CCg, nCubeg, SortPartg, OuterReigong,BeginEndCubeg, RepForceg);      
      
      cudaError_t error12 = cudaGetLastError();
printf("CUDA error12: %s\n", cudaGetErrorString(error12));
      CHECK(cudaDeviceSynchronize());
      cudaError_t error26 = cudaGetLastError();
printf("CUDA error26: %s\n", cudaGetErrorString(error26));
   }

不妨发一下你的kernel的代码,然后在kernel后面立刻取回(cudaXXXSynchronize()这里可以省略,如果你后面立刻同步取回的话),而不要在另外一个函数里,避免你的CPU上有代码中间干别的事情,而对你造成了错觉。然后我看下。

11
CUDA / Re: Invalid __global__ write of size 4
« 于: 十一月 02, 2022, 12:40:16 pm »
/data/apps/cuda/11.4/bin/nvcc -G  --device-c system.cu -o system.o    请问是类似于这样对吗

如果你需要分步编译的话,则是这样的。需要-G和-dc的。没错。(或者你至少使用-dc和-lineinfo也可以,如果-G
编译出来的代码,运行速度极度缓慢,而在cuda-memcheck下半天没结果的话)。

12
CUDA / Re: 从Device端传回Host端的时候cudaMemcpy函数不起作用
« 于: 十一月 02, 2022, 12:38:17 pm »
这种有可能的,最常见的情况是kernel启动配置失败(没有执行),而你没有使用经典的2次检查的方式来捕获错误,也就是是根据在Kernel启动<<<>>>后面的立刻一次CHECK(cudaGetLastError()); 和在后续的最近一次地方的同步调用(例如对cudaXXXSynchronize()或者其他同步操作的CHECK())。这两个分别代表检测"kernel是否能启动", 和"kernel启动后,后续执行是否成功(例如没有因访存而挂掉).

常见的是漏掉前者,而前者的错误可能会被后续的其他操作的某个地方的成功给覆盖掉。而只有后者的错误才是持续能返回的。所以往往常见认为kernel已经执行好了,结果没动弹,往往漏在这里。

当然,不能排除其他方面的原因,例如是否你的kernel真的啥都没干,不过前面说的是比较常见的情况,可以先排查下。

以及,不妨就在"这个kernel"中使用printf,看看能否发现值已经被修改了。注意不是在另外一个kernel中,也不是观察"原始输入值", 这样可以很大程度的避免一些幻觉。

13
CUDA / Re: 从Device端传回Host端的时候cudaMemcpy函数不起作用
« 于: 十一月 02, 2022, 12:35:15 pm »
请问一下我在GPU上计算完结果之后,想要传回CPU端,写了如下代码
void CSPH::ParticlesDataReturn()
{
   CHECK(cudaMemcpy(TotalParCoor, d_TotalCoor, sizeof(double3)* nTotalPar, cudaMemcpyDeviceToHost));
   CHECK(cudaMemcpy( Press,d_Press, sizeof(float)*nTotalPar, cudaMemcpyDeviceToHost));
   CHECK(cudaMemcpy(velrhop,d_Velrhop , sizeof(float4)*nTotalPar, cudaMemcpyDeviceToHost));
   CHECK(cudaMemcpy(SortPart, d_SortPart, sizeof(int)*nTotalPar, cudaMemcpyDeviceToHost));
   CHECK(cudaMemcpy(RepForce, d_RepForce, sizeof(float3)*nFluidPar, cudaMemcpyDeviceToHost));
   cudaError_t error = cudaGetLastError();
printf("CUDA error: %s\n", cudaGetErrorString(error));
   
for(int i=0;i<nFluidPar;i++)
{
   if(SortPart==0)
printf("%d, %f, %f, %f aaaaa\n ",i,RepForce[SortPart].x,RepForce[SortPart].y,RepForce[SortPart].z);

}
}
在传回的这些数据中心,只有RepForce[].x,RepForce[].y,RepForce[].z的值还跟程序刚开始在CPU端赋的初值相同,而其他数组的值输出来就是在GPU上计算完成后的值。我在核函数中输出了d_RepForce[]的值是对的,感觉是从GPU传回CPU的时候cudaMemcpy函数没起作用,并且用cudaGetLastError();也显示no error。请问这是怎么回事呀

这种有可能的,最常见的情况是kernel启动配置失败(没有执行),而你没有使用经典的2次检查的方式来捕获错误,也就是是根据在Kernel启动<<<>>>后面的立刻一次CHECK(cudaGetLastError()); 和在后续的最近一次地方的同步调用(例如对cudaXXXSynchronize()或者其他同步操作的CHECK())。这两个分别代表检测"kernel是否能启动", 和"kernel启动后,后续执行是否成功(例如没有因访存而挂掉).

常见的是漏掉前者,而前者的错误可能会被后续的其他操作的某个地方的成功给覆盖掉。而只有后者的错误才是持续能返回的。所以往往常见认为kernel已经执行好了,结果没动弹,往往漏在这里。

当然,不能排除其他方面的原因,例如是否你的kernel真的啥都没干,不过前面说的是比较常见的情况,可以先排查下。

14
深度学习讨论 / Re: 线程束分化
« 于: 十一月 02, 2022, 12:28:29 pm »
想请教一下:N卡上的线程束分化,看cuda c编程权威指南里面提到,当出现线程束分化的时候,同一个线程束中,满足con的线程束会执行if指令,不满足的会等待,我想问的是,线程束在指令调度单元下,在一个周期内不都是执行的相同的指令吗,为何不满足的是等待呢?按自己的理解应该是不满足的也会执行该指令,只是该结果被某种GPU策略给丢弃了。不知道哪里理解的不对?谢谢

当一个warp的内部出现divergent branch的时候,例如你用if(cond){body1} else {body2}的方式,那么在较新的卡上,支持独立线程调度的时候,有可能warp中的满足cond的线程们在执行条件body1的时候,剩下的线程们并不等待,也可能会执行body2的。

其次就是,如果body1和body2的跳转没有被编译成分支,而是带有@p body1的这样的指令序列的时候,则对于warp中的所有线程都会执行body1,但是不满足p为真的那些线程,也就是类似@false XXXX,这里的XXXX的执行效果会被丢弃。(丢弃是指,普通计算不写入结果;访存指令不evaluate地址,不执行实际上的存储器访问。所以是安全的。)

所以你之前的理解在某种程度上是对的。

15
CUDA / Re: 使用了共享内存最终时间没有没有减少
« 于: 十一月 02, 2022, 12:23:24 pm »
那意思就是如果想要借助共享内存来减少存取延时,必须要在一个核函数范围内了吗? 
假设想要把这3个kernel合并在一个kernel中,因为寄存器的原因导致launch失败【寄存器数量超出硬件限制】,即不能合并成一个kernel,那这种情况下只能把 U数组 存在global memory了吗?

(1)首先说,使用global memory并不一定会慢的(或者说使用shared memory一定会快是错觉)。
(2)其次是,如果你非要这样做的话也可以。同时计算能力5.0以上的硬件(最近8年)并不存在"寄存器超出硬件数量限制的问题", 因为编译器最多会为5.0+的卡生成使用最多255个常规寄存器的代码,而5.0+的硬件本身最多一个线程也可以分配255个寄存器(但此时可能最大的block形状受到限制,例如256个线程的block)。当然,你如果还在使用3.X的老卡,那就没办法了. :(

页: [1] 2 3 ... 45