[attach]3444[/attach]
谢谢指正。可是，如图所示，这是CUDA自带的原版归约优化例子中的第一次优化程序，红色字体说% “is very slow”. 这句话应该如何理解呢？
人家说的是：
（1）warp内严重分支会影响效率
（2）对2 * s求余会严重影响效率
但人家没说对2求余会严重影响效率！！！！！！！
知道了，谢谢。
请问Exception里的outofrangeload是什么错误？提示access violations的地方就是OutOfRangLoad.
LZ您好：
access violations一般是访存越界。
而OutOfRangLoad值得应该就是本意，告诉您在读取的时候越界。
祝您好运~
今天写完小波去噪程序，但发现很慢。要9秒多的时间。虽然数据量有点大，但也不至于慢到这程序吧。
不知道瓶颈在哪里，版主能否给出一些优化建议。或者给出更合理的算法。
#include<stdio.h>
#include<math.h>
#include <stdlib.h>
#define m 160000
#define n 128
#define nmi 7
#define scale 6 //scale分解级数，
#define wlen 6 //wlen小波长度
float c0[m][n];
__global__ void denoise(float *cd)
{
int tid=threadIdx.x;
int xr=threadIdx.x+blockIdx.x*blockDim.x;
int mid;
float p,q;
int sca[nmi]={128,64,32,16,8,4,2};
int flg[nmi]={0,128,192,224,240,248,252};
float h[wlen]={0.332671,0.806892,0.459878,-0.135011,-0.085441,0.035226}; //小波系数
float g[wlen]={0.035226,0.085441,-0.135011,-0.459878,0.806892,-0.33267}; //尺度
__shared__ float appx[n*2],detl[n];
while (xr<m*n)
{
appx[tid]=cd[xr];
__syncthreads();
/************** DWT *************/
for (int j=1;j<=scale;j++)
{
if (tid<sca[j-1] && tid%2==0)
{
p=0;
q=0;
for (int k=0;k<wlen;k++)
{
mid=tid+k;
if (mid>=sca[j-1])
mid-=sca[j-1];
p+=h[k]*appx[flg[j-1]+mid];
q+=g[k]*appx[flg[j-1]+mid];
}
appx[flg[j]+tid/2]=p;
detl[flg[j]-n+tid/2]=q;
}
}
/************** FILTER *************/
if (tid<sca[1]) {detl[flg[1]-n+tid]=0;}
if (tid<sca[6]) {appx[flg[6]+tid]=0;}
/************** IDWT *************/
for (int k=scale;k>0;k--)
{
if (tid<sca[k])
{
p=0;
q=0;
for (int s=0;s<wlen/2;s++)
{
mid=tid-s;
if (mid<0)
mid+=sca[k];
p+=h[2*s]*appx[flg[k]+mid]+g[2*s]*detl[flg[k]-n+mid];
q+=h[2*s+1]*appx[flg[k]+mid]+g[2*s+1]*detl[flg[k]-n+mid];
}
appx[flg[k-1]+2*tid]=p;
appx[flg[k-1]+2*tid+1]=q;
}
__syncthreads();
}
cd[xr]=appx[tid];
__syncthreads();
xr += blockDim.x*gridDim.x;
}
}
int main()
{
for (int i=0;i<m-1;i++)
for (int j=0;j<n;j++)
c0[i][j]=i+0.5*j;
float *cd;
cudaMalloc(&cd,m*n*sizeof(float));
cudaMemcpy(cd,c0,m*n*sizeof(float),cudaMemcpyHostToDevice);
denoise<<<3,n>>>(cd);
cudaMemcpy(c0,cd,m*n*sizeof(float),cudaMemcpyDeviceToHost);
for (int i=0;i<n;i++)
printf("%.3f\t",c0[0][i]);
}
LZ您好：
新问题请您开新帖讨论，您这不到30楼已经换过数个主题了。
以及，看了您的代码，简单说一下：
因为我不懂小波变换，也无法根据您的代码立即自学学会该算法的原理，并想出一个更合理的算法，因此，将不就该算法自身以及是否在算法层面能有改进进行讨论。如其他网友/斑竹/总版主/NV支持人员有您的同行，您可以和他们就小波变换算法进一步深入讨论。
针对您的具体算法，大致说一些明显的问题
1：您的线程数量太少，远远少于一般的习惯用量，而您的计算数据量并不算少。可见您的线程规模偏小并不是因为计算规模偏小造成的。如果在您的算法许可内增加计算用的线程数量，那么将有助于提高效率。（请注意，您的算法是否允许进一步拆分以增加线程数量，这一点我无法判断，请您自行考虑）
2：部分对shared memory写入的操作并没有进行同步，请您注意并验证这里。
3：您的代码中有非warp对齐的if判断，并且在某些循环中存在越循环干活的线程数量越少的情况，如果您能在算法等价的前提下，修改实现，解决这个问题，将有助于提高效率。
大致如上，您也许需要在算法等价的前提下，修改实现使得适合GPU计算，并且启动足够多的线程才能达到较好的效果。
此外，您可以跑visual profiler，通过参考其统计结果来进行针对性的优化。您的程序逻辑较为复杂，斑竹实无法人肉profiler判断。
祝您编码顺利~
1、算法的前半部分为卷积运算，越到最后参与卷积的线程越少，直到只有两个线程。后半部分刚好相反，从两个线程开始，依次翻倍增加。类似于归约运算，这种算法中的线程浪费似乎无法避免。
2、有一个地方我以为你会提到，我直接问吧。我的核函数里有四个现成的一维数组，每个活动的线程都会用到这四个数组。如果我像上面这样写，不加任何前缀，这几个数组是放在GLOBAL里的吗？如果我想把它放在共享存储器里面，理论上速度会不会有提升？我做过实验，感觉没什么提升，但我不能排除是否因为我的程序不合理。
伴随而来的一个问题，直接在共享存储器里定义数组，又不允许初始化，我试过一个一个单元来赋值，感觉很烦。一般是怎样对共享内存里的数组进行赋值呢？
每个线程有单独的寄存器，如果一个BLOCK里的很多线程是合作的关系，使用共享内存里的数据，那么每个SP的私有寄存器用来干什么？
3、我的GPU只有384核，而且一次处理的数据是128个，您提到的第1点我暂时没有相应的算法。
4、至于您提到的第二点，我原本在其它位置也有一些同步指令，但是发现不要那些同步指令也能得出正确的结果。令我难以理解的是，有些地方加了同步指令反而只能得到一部分正确结果。。。我想问同步指令除了影响线程同步之外，还会不会对效率产生影响，比如有的线程跑得很前，有的很落后，在很远的地方再进行同步，这样对效率有影响吗？
5、我对CUDA还有很多知识漏洞，很多分析工具也还不能充分利用，profiler我要先学习学习才能用。
6、我是根据CPU版本来改写的这个程序，CPU跑出来是620ms。我对这个GPU程序都绝望了。
7、谢谢您。
LZ您好：
1：我不懂卷积运算，也不懂卷积运算的高效实现，请您根据您的实际问题决定，为避免误导，我不能发表意见。
2.1：像您这样这样把数组写在kernel里面，是无论如何不会放置到global memory中的。一个在kernel中的普通的数组，如果在编译的时候起下标访问规律已知，那么将会使用寄存器来保存（在寄存器数量不够的时候，编译器将会选择最常用的保存在寄存器中，不常用的spill到local memory）；如果在编译的时候下标访问规律未知，那么数组将整体被放置到local memory中。
以及，如果是像您的这种完全常数的小型数组，如果编译时下标访问规律已知，还有可能编译以后直接作为立即数或者存放入constant cache中。
2.2：使用shared memory是否能提高性能，这个不一定。如果之前已经被编译成了立即数访问，那么已经是最快的了。如果之前是constant cache访问，那么在warp内访问行为不一致的时候，shared memory可能会更快一些。以及shared memory要注意避免 bank conflict。
2.3：shared memory一般使用一组线程写入进行初始化。
2.4：每个线程的私有的SP干剩下的绝大多数的事情。一般在kernel中定义的变量都会放入寄存器，定义的数组会尽量放入寄存器（需要下标访问规律在编译时可知），其他中间变量也会使用寄存器。基本上所有的计算都只能和寄存器直接打交道。
3：您的理解是错误的，时间原因，这里先不展开了。通过您的这句话，可以看出您对基本概念的理解并不到位，以及您对您的所需要完成的任务的理解也不到位。
4：请通过适当的同步安全合理地使用shared memory，而不要因为您自己写错而怀疑这个怀疑那个。
以及，同步指令确实有一定开销，但是不以正确和可靠为前提，速度有什么意义呢？
5：请尽快学习profiler，这是一个很有用的工具。
6：火车刚造出来的时候也没马车快，现在呢？
7：不客气。
祝您好运~
补充一下，一般CPU的算法需要在算法等价的前提下对具体实现充分进行针对GPU特性的修改，这样才能发挥出GPU的威力，或者说体现出程序员的价值。
直接套用，或者简单修改，常常无法发挥GPU的性能。
我正在参考sample里面的一维haar小波变换，但是例子里每个block用了512个线程，我的电脑运行不了。无法跟踪数据，很难看出程序的思路。
有两个指令看不明白
[attach]3450[/attach]
定义的data0,data1是float型变量，而等号后面的是一个共享存储器空间，把一个空间里那么多数据赋值给一个变量，这怎么理解呢？
楼主您好，
如果有类似:
float data = set[n];
的代码，则是将set中的第n个元素赋值给data变量。而不是您原本认为的“将一堆元素给一个变量”，希望您能理解。
感谢来访。
LZ您好：
1：我无法为您解释CUDA Samples中的小波变换的例子，我不懂小波变换，CUDA Samples根据授权协议，也不享受任何支持的。
2：任何计算能力版本的NVIDIA GPU都至少支持每个block内部有512线程的，因而您不能用这个理由来说明您的问题，我将视之为某种搪塞或者直接无视此句。
3：原先的写法是正确的，您的问题我无法理解。请问您真的是C语言的程序员么？（以下省略若干字）
祝您好运~
tid=threadsIdx.x;
N个线程并行运算，tid不就是同时为0～N吗，data至少有N个float单元吧。
我没有说过我是程序员，我只是校园里二级C水平的入门菜鸟。如果请教的问题过于简单，请原谅。
N个线程里有N个，你可以当作一个线程只有一个。
而一个线程里只有一个确定的下标，从而有一个确定的值，是不是你就容易理解了。
那么说DATA是每个线程的私有寄存器？而且每个私有寄存器里都有一个同名为data0的变量？
LZ您好：
写程序的人都是程序员，无论是专业的还是业余的，无论是资深的还是入门的。
之所以问你是不是c语言的程序员，用以在于，在其他的一些语言中（比如fortran），可以直接一个矩阵对一个矩阵进行运算，而不用像C语言里面那样拆开控制元素运算。
如果您之前是fortran程序员，可能会受到这个习惯的影响，从而产生错误的理解或者一时的误解。因而提醒您注意一下。
此句的具体原因，横扫斑竹已经详细解释了，不再赘述。
祝您好运~
LZ您好：
在kernel里面定义的变量和数组都是每个线程私有的，除了__shraed__ 声明的，这个是整个block共有的。
您可以将“一个线程”简单看做“一个十分微小但是全功能的计算机”，他们之间用线程编号区分，以及每个“计算机”自己会跑一份kernel代码，那么自己代码中定义的变量当然是私有的，和其他“计算机”无关。
大致如此，祝您好运~