最新帖子

页: [1] 2 3 ... 10
1
CUDA / 改造成kernel时,参数很多,怎么办?
« 最后发表 作者 master 八月 30, 2023, 09:57:26 pm »
想把原来在CPU下运行的子程序,改造成CUDA FORTRAN里的kernel。遇到一些问题,想请教大侠指点下。

原来的程序把变量定义在一个module里,其中还用save定义了一些全局变量。想要改造的这个子程序中用到了不少这些全局变量,有几十个。这个子程序开头use variables来使用这些全局变量,varialbes是存放这些变量的module。

这样的话,改造前这个子程序的参列表里就不需要把这些全局变量也列上,只列几个局部变量即可。

现在把子程序改造成kernel时就有点麻烦。这些全局变量是由CPU先计算好,然后调用kernel,这个kernel只是用了一下这些变量值,并不改变它们的值。

现在的改造思路是:把这些全局变量先传递到device上的相应数组,然后调用这个kernel,那么这个kernel的参数列表就将会很长(即需要把这些变量都写进去),因为kernel里不能定义全局变量。

这些全局变量在CPU里已经计算好,不再变化,调用kernel也不会导致它们变化。

kernel需要被反复调用数百次,每次调用时都有这些参数,感觉一是这个kernel后面括号的参数列表太长,也不好看。还有就是每次调用kernel都有这些参数,会不会影响效率。

不知道有没有遇到这种问题的,是如何解决的呢?遇到这种情况,该如何改造kernel?有没有办法使参数里面不写这些变量呢?
2
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即可。
3
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, 只是套上了图形化的界面了),十分方便,欢迎尝试。

4
CUDA / Re: 多个线程进行判断修改同一个全局内存遇到的问题
« 最后发表 作者 屠戮人神 七月 31, 2023, 10:30:15 pm »
首先说,你这试图做的是常见的规约运算,对多个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()),你随意即可。
5
CUDA / Re: 多个线程进行判断修改同一个全局内存遇到的问题
« 最后发表 作者 屠戮人神 七月 31, 2023, 10:26:11 pm »
程序代码: [选择]
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的,因此无需繁琐的人力调试时间,快速面对老板和市场的挑战。

6
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了。
7
CUDA / Re: 共享内存数组大小及bank conflict
« 最后发表 作者 屠戮人神 七月 31, 2023, 08:53:12 pm »
在第一个问题里面,你也可以选择用cuda-gdb或者nsight来定位到越界的地方,从而减少资源分配到正好足够。也是可以的。但是这样对于开发者的时间不一定合算。你自己决定要不要修复BUG,还是选择掩盖BUG。
8
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。
9
CUDA / cuda-gdb 显示的错误信息
« 最后发表 作者 2017012835 五月 15, 2023, 09:16:50 am »
请问一下,用 cuda-gdb 调试程序输入了 r 命令,得到了图片上的报错内容,之前的报错都是定位到了程序中的某一行,但现在这个报错不知道说的是什么意思,请问出现这样的错误可能是哪里的问题呢?
10
CUDA / 多个线程进行判断修改同一个全局内存遇到的问题
« 最后发表 作者 jinyer 三月 30, 2023, 11:21:58 am »
程序代码: [选择]
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吗?
页: [1] 2 3 ... 10