1

CUDA / Re: 共享内存数组大小及bank conflict

« 于: 三月 01, 2023, 05:40:09 pm »

您好，您第一个问题是不是数组索引越界了，for循环里的S_A[threadIdx.x * 12 + j]，其中threadIdx.x最大是127，当它取最大值（或没有取到最大值时），threadIdx.x * 12 + j都超过了您设置的S_A[128]里面的128

您好，我按照您说的索引越界问题进行了修改将S_A定义设置为__shared__ float S_A[1536]，程序可以正确运行，但是将定义设置为__shared__ float S_A[1280]程序也能正确运行，所以对共享内存数组大小问题还是有疑问。


另外，在我的核函数中对共享内存的使用如下

程序代码: [选择]

kernel2(float* A,float* M, float* Y,float* X)
{
    const int n = blockDim.x * blockIdx.x + threadIdx.x; 
       E[n]=0;

        __shared__ float S_E[sizeof(float) * 320];       
        __shared__ float S_A[sizeof(float) * 320];   

       S_E[threadIdx.x] = E[n];
__syncthreads();
        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];     
atomicAdd(&S_E[threadIdx.x],S_A[threadIdx.x * 12 + j]/2); 
        }
__syncthreads();


}请问如何才能最大程度减轻共享内存S_A和S_E共享内存的bank冲突问题？

2

CUDA / 共享内存数组大小及bank conflict

« 于: 二月 28, 2023, 06:04:33 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冲突？

3

CUDA / Re: 关于共享内存的使用

« 于: 十二月 07, 2022, 10:03:45 am »

你这样直接初始化了shared memory里的内容，然后大家集体就这样用，显然不行的。你需要添加同步的，在每次对shared memory的改动后，和下次读取前。

此外，同步你还不能放在if里面，你可能需要处理一下你的代码架构。

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

我按照您说的问题进行了修改，现在程序可以正确运行，但是耗时却比使用全局内存还多，请问使用共享内存不一定会提高程序的运行速度吗？

4

CUDA / 关于共享内存的使用

« 于: 十一月 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];
    }

   }
}

5

CUDA / 将循环移至核函数内计算出错

« 于: 十一月 18, 2022, 11:57:07 am »

以第一种方式进行循环时可以得到正确的计算结果，但是以第二种方式循环时计算出现错误，请问是什么原因导致计算错误？

6

CUDA / 用nsight compute分析程序时报错

« 于: 十一月 14, 2022, 11:20:40 am »

用nsight compute 分析具体核函数时报错No report sections available.See Profile options，应该怎么解决？

7

CUDA / Re: .cu文件用nvcc编译为.exe后不能运行

« 于: 十一月 03, 2022, 02:50:19 pm »

如同之前说的，如果单独运行exe打开文件失败，在VS下运行却可以打开文件，最大的可能是路径问题。你可以直接使用绝对路径，来避免这种情况。例如说："d:\\files\\a.txt"， 类似这种路径。

如果你还有能力，可以考虑在Windows下使用GetModuleFileNameA(), 使用NULL作为句柄，这会得到当前exe的全路径，然后你可以从全路径出发，从最后一个反斜杠后面开始替换，例如：
"C:\\dog\\cat\\wolf.exe"
和同时在C盘的dog文件夹的cat文件夹下面，和exe在一起放置有data.txt
则你可以从最后一处反斜杠开始替换，得到
"C:\\dog\\cat\\data.txt"
这样就不用写绝对路径了，灵活一点。

最后发现是主机端读取一个txt文本中的二维数组时出现了问题，已经解决，感谢您提供的思路

8

CUDA / Re: .cu文件用nvcc编译为.exe后不能运行

« 于: 十一月 01, 2022, 10:14:24 am »

此外，你的程序如果在VS下运行，可能默认当前路径不同，如果你用了一些数据文件之类的，可能在VS下能打开，直接exe打不开从而导致运行失败之类的，都是有可能的。不妨排查一下。至于NSight性能调优，你先解决了你的运行问题再说吧。不要一个帖子同时说两点，这样可能会弄乱你对问题的所在的分析的。

您好，经过测试，exe不能运行的原因应该是我读取了一些txt文本中的数据，在VS下能够读取并运行程序，但是exe打不开所以运行失败，不知道这个问题应该怎么进行解决？

9

CUDA / Re: .cu文件用nvcc编译为.exe后不能运行

« 于: 十月 23, 2022, 09:27:22 pm »

我想用nsight compute分析一下程序，但是发现程序用，nvcc编译为exe后运行不了，并且在nsight compute选择该exe文件launch后显示连接失败。但是程序在VS2019上可以成功运行，不知道是什么原因。

10

CUDA / .cu文件用nvcc编译为.exe后不能运行

« 于: 十月 19, 2022, 09:09:47 pm »

.cu文件在visual studio2019上可以正常运行，但是在命令行中用nvcc kernel.cu -o kernel编译生成exe后运行exe程序不能运行。