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吗？