列出帖子

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


显示所有帖子 - 屠戮人神

页: [1] 2 3 ... 46
1
CUDA / Re: cuda-gdb 显示的错误信息
« 于: 七月 31, 2023, 10:42:17 pm »
你好,根据图片,你挂在了__ldg()函数的内部,八成是你使用了/计算了错误的地址/下标/索引之类的缘故。

你可以使用bt/info stack之类,进行回溯/查看调用堆栈(倒过来的)。
例如:
) info stack
#0 ... at sm_XX_intrinsic.hpp:134
#1 ... at your.cu:2333
#2 ... at your2.cu:6666
查询后面数字较大的frame所在的文件和行号即可,后面的一般的都是你自己的。这样即可快速定位。

此外,需要说明的是,现在visual studio code在linux上可以作为cuda-gdb的GUI外壳进行调试,直接从VS Code的图形化界面里查看/调试,无需记忆cuda-gdb命令(底层依然是cuda-gdb, 只是套上了图形化的界面了),十分方便,欢迎尝试。

你也可以临时性的将__ldg给重新define成普通的指针访存,这样可以直接cuda-gdb在top顶层frame显示你自己的代码的行号,而无需查看堆栈回溯信息,可能有的时候更方便一遍。调试完成了,再取消掉define即可。

2
CUDA / Re: cuda-gdb 显示的错误信息
« 于: 七月 31, 2023, 10:40:00 pm »
请问一下,用 cuda-gdb 调试程序输入了 r 命令,得到了图片上的报错内容,之前的报错都是定位到了程序中的某一行,但现在这个报错不知道说的是什么意思,请问出现这样的错误可能是哪里的问题呢?

你好,根据图片,你挂在了__ldg()函数的内部,八成是你使用了/计算了错误的地址/下标/索引之类的缘故。

你可以使用bt/info stack之类,进行回溯/查看调用堆栈(倒过来的)。
例如:
) info stack
#0 ... at sm_XX_intrinsic.hpp:134
#1 ... at your.cu:2333
#2 ... at your2.cu:6666
查询后面数字较大的frame所在的文件和行号即可,后面的一般的都是你自己的。这样即可快速定位。

此外,需要说明的是,现在visual studio code在linux上可以作为cuda-gdb的GUI外壳进行调试,直接从VS Code的图形化界面里查看/调试,无需记忆cuda-gdb命令(底层依然是cuda-gdb, 只是套上了图形化的界面了),十分方便,欢迎尝试。


3
首先说,你这试图做的是常见的规约运算,对多个predicate或者说布尔值,进行or规约,最后形成标量。即:
bool final_result = (checkResult[0] != 0) || (checkResult[1] != 0) || (checkResult[2] != 0) ...
final_result = !final_result;

你的代码写法是常见做法,无需怀疑。但是至少这里有错误:
cudaMemset(..., 1, ...);
这里将会将你的初始化设定为0x01010101这么一个16进制数, 而非你想象中的0x00000001;
前者转换成10进制,正好是16843009。

所以你的kernel就变成了,如果有任一一个检测到的值非0,则结果清零。否则结果为16843009。

而你观察到的有线程写入了0,最后cudaMemcpy回来,却依然非零,那只是你的幻觉(很常见,很容易多次调试的值你给弄混,所以我一般调试总是做笔记)。

这个世界是科学的,考虑到这里,我建议你直接对常见运算(布尔值规约),使用成品thrust库进行。即:
存在int *check为设备上的缓冲区,有N个元素,则:

#include <thrust/logical.h>
#include <thrust/functional.h>
#include <thrust/execution_policy.h>
...
int main()
{
     bool result = thrust::any_of(thrust::device, check, check + N, [] __device__(int i)
    {
        return i != 0;
    });
    result = !result; //你的要求是反的
    ...
}

这样只需要提供你的核心逻辑(元素i != 0), thrust配合C++将自动生成kernel,里面填充上你的核心语句(非0判断),然后将得到的结果自动规约(any_of). 而thrust是production ready的,因此无需繁琐的人力调试时间,快速面对老板和市场的挑战。

此外,thrust会尝试尽量短路操作,如果中途规约(布尔值的or运算)的时候,他如果中途发现了任意1个值已经为true了,则后续的值的判断可能不再进行,将可能不再回将所有数据都处理一遍,比我们想象的一般要聪明一点。

此外,该操作需要开启nvcc的extended lambda支持,手册里有,我建议你阅读一下。(不开启你需要重载一个operator()),你随意即可。

4
程序代码: [选择]
int *d_flag;
checkCudaErrors(cudaMalloc((void**)&d_flag, 1*sizeof(int)));
cudaMemset(d_flag, 1, sizeof(int)*1);  //初始化为1
block = 96;
isIter<<<(M+block-1)/block,block>>>(d_flag, d_check);
isIter核函数想实现的功能是 :flag[0]初始化是1,若check数组有一个数不为0,则将flag[0]置为0
程序代码: [选择]
__global__ void isIter(int *flag, int *checkResult) {
int idx = threadIdx.x + blockDim.x * blockIdx.x;
if(idx < M){
if(checkResult[idx] != 0){
flag[0] = 0;
}
}
}
可这样老是会输出flag[0]=16843009(也有部分线程进入if语句里边,输出flag[0]=0),可将flag[0]复制到CPU读取的时候是16843009,请问这是为什么啊?
flag存储在全局内存,不应该若有线程修改的话值就变为0,没有线程修改还是保持原值1吗?

首先说,你这试图做的是常见的规约运算,对多个predicate或者说布尔值,进行or规约,最后形成标量。即:
bool final_result = (checkResult[0] != 0) || (checkResult[1] != 0) || (checkResult[2] != 0) ...
final_result = !final_result;

你的代码写法是常见做法,无需怀疑。但是至少这里有错误:
cudaMemset(..., 1, ...);
这里将会将你的初始化设定为0x01010101这么一个16进制数, 而非你想象中的0x00000001;
前者转换成10进制,正好是16843009。

所以你的kernel就变成了,如果有任一一个检测到的值非0,则结果清零。否则结果为16843009。

而你观察到的有线程写入了0,最后cudaMemcpy回来,却依然非零,那只是你的幻觉(很常见,很容易多次调试的值你给弄混,所以我一般调试总是做笔记)。

这个世界是科学的,考虑到这里,我建议你直接对常见运算(布尔值规约),使用成品thrust库进行。即:
存在int *check为设备上的缓冲区,有N个元素,则:

#include <thrust/logical.h>
#include <thrust/functional.h>
#include <thrust/execution_policy.h>
...
int main()
{
     bool result = thrust::any_of(thrust::device, check, check + N, [] __device__(int i)
    {
        return i != 0;
    });
    result = !result; //你的要求是反的
    ...
}

这样只需要提供你的核心逻辑(元素i != 0), thrust配合C++将自动生成kernel,里面填充上你的核心语句(非0判断),然后将得到的结果自动规约(any_of). 而thrust是production ready的,因此无需繁琐的人力调试时间,快速面对老板和市场的挑战。


5
CUDA / Re: 有关对缓存优化的问题
« 于: 七月 31, 2023, 09:04:11 pm »
关于你的两个问题,

1)有多方面的资料。一般的来说,对于local memory访问,因为总是存在强制性自动交错的。即warp中的32个线程,每个人拿到的相同的1个4B变量地址(local memory中),实际上的地址会被自动交错到连续的128B上。所以一般来说local memory你无需担忧。可以认为是正常情况下总是100%合并的L1 - L2 - 显存的访问(在你的计算能力架构是上)。

而对于global memory的访问,在你的计算能力架构上,读取是会通过L1 <- L2 <- 显存或者内存的。而写入则一般会L1 -> L2, 然后再不定的时候L2 -> 显存/内存。你搜索nvidia + gpu + l1 write through + l2 write back,可以获取更多信息。

2) 你的第二个问题,虽然你询问的是L1的hit ratio,这个不一定是越高越好,要看问题(例如streaming风格的访问,例如常见的2个向量的加法,直接回写显存)。但是profiler给出的另外一个方面是一个重大问题,它指出(你没有源代码)你的global memory的load和store,均是warp里的每个线程地址分散的很厉害的情况。虽然看不到代码,但是建议先解决这个问题,再看时间和cache命中率的变化。可能你解决了这个问题,你发现你不在需要解决问题2了。

6
CUDA / Re: 共享内存数组大小及bank conflict
« 于: 七月 31, 2023, 08:53:12 pm »
在第一个问题里面,你也可以选择用cuda-gdb或者nsight来定位到越界的地方,从而减少资源分配到正好足够。也是可以的。但是这样对于开发者的时间不一定合算。你自己决定要不要修复BUG,还是选择掩盖BUG。

7
CUDA / Re: 共享内存数组大小及bank conflict
« 于: 七月 31, 2023, 08:51:19 pm »
1:我在核函数中设置线程块的大小为128,我需要将float类型的变量定义为共享内存变量,发现当将共享内存数组定义写成数组长度__shared__ float S_A[128]时程序运行是错误的,当把定义写为__shared__ float S_A[sizeof(float)*320]时可以得到正确的结果。共享内存不是可以直接定义数组长度吗?
2:我在核函数中共享内存使用如下
程序代码: [选择]
kernel2(float* A,float* M, float* Y,float* X)
{
    const int n = blockDim.x * blockIdx.x + threadIdx.x; 
     
        __shared__ float S_A[sizeof(float) * 320];       
           
        for (int j = 0; j < 12; j++)
        {           
                S_A[threadIdx.x * 12 + j] = A[n * 12 + j];

                S_A[threadIdx.x * 12 + j]=M[Y[n]]+X[n * 12 + j];           
        }
__syncthreads();


}

因该怎么修改程序消除或减轻共享内存的bank冲突?

这是两个独立的问题:

第一个问题是常见现象, 如果正好分配(按照元素数量) X[N]个,结果却发现kernel跪了, 而选择其他的无关的X[N + K], 或者X[N * K]的形式,扩大的容量,结果却能跑出来。 这往往代表了你下标越界。

这种情况有两种处理方式,一种就是你选择的扩大分配容量,来掩盖掉越界的BUG(BUG依然存在,只是可能不触发了),这样如果资源足够的情况下,完全可以的。BUG不一定需要立刻修复,只要代码能跑就行。留给以后的同事也算是一段因缘。

第二个问题,则是说,对于每个线程(假设threadIdx.x不重复你这里)的形如s_A[threadIdx.x * 12 + j], where j ∈ [0, 12), 的访问, 或者用中文说,每个线程要留出12个4B元素的空间,这种每次warp整体的访存,的确会导致bank conflict。根据质因数分解我们可以知道scale = 12 = 3 * 4, 这样会导致4-way的bank conflict。你可以从profiler直接得到4-way bank conflict的信息, 或者也可以直接例如例如说warp里的lane 0,8,16,24均会访问同一个bank。

想解决这个问题,一般采用最近的不含有质因数2的scale值。例如你可以从12变成13,即s_A[threadIdx.x * 13 + j], j依然∈ [0, 12),这样每个线程分配了13个元素的空间,而不在是12个。此时对于warp的任意一次访存, x * 13 != x' * 13 (mod 32)均成立, 即将不再存在bank conflict。

8
CUDA / Re: 请教有关FP32和FP64相关问题
« 于: 二月 13, 2023, 04:08:56 pm »
为什么我在代码中 所有的计算部分已经全都是用的float类型(可能一部分是FP64单元去计算了),如Nsight Compute分析所示,为什么还提示我让我用FP32,如何让FP64不去参加运算?在这种情况下请问怎么进行优化呢?

关于你的两个问题:
(1)为何报告FP64 Pipe几乎用满了,但是却"所有的计算代码已经全部使用的是float类型"的问题,可能有这几个原因:
表达式中有无不小心引入的double变量?或者有无不小心引入的double<->int, double<->float之间的转换?以及,特别检查有无不含f结尾的常数,例如说1.23而不是1.23f,前者会引入自动的到double的promotion转换和计算代价。

(2)为何报告大比例的FP32不能融合成a * b + c的方式,而是单独的a * b, 和x + c?这个考虑到(1)问题的存在,而现在报告FP32 Pipe才3%的使用率,不妨先解决了(1)然后再看看本情况的新报告,和新变化。


9
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可以被选中+发射指令的。

10
CUDA / Re: 关于共享内存的使用
« 于: 十二月 15, 2022, 02:03:32 pm »
我按照您说的问题进行了修改,现在程序可以正确运行,但是耗时却比使用全局内存还多,请问使用共享内存不一定会提高程序的运行速度吗?

使用共享内存不一定会总是提速的,这是一个正常现象。
不是特殊的银弹, 本身不具有灭杀狼人的魔法效果。

11
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参与了下标计算之类的式子的话)。

12
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里面,你可能需要处理一下你的代码架构。

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


13
您好,我按您说的检查了,但是发现您说的这三种情况都没有发生。我用cuda-memcheck检查了内存。总是显示是这个函数上面核函数非法访问内存的问题,请问上面的核函数出问题了,是否会造成这个函数中输出错误的这种情况呢?

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

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

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

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

15
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都用。

页: [1] 2 3 ... 46