非常感谢！！！
int num_ees = 12;
int num_step = 16;
int num_gg = 10241024;
int num_pp = 10248;
int num_dims= 1024*2;
int num_segments= num_dims/num_step;
cudaMalloc((void **) &d_data_gg, num_ees * num_step * num_gg * sizeof(char) );
cudaMalloc((void **) &d_data_pp, num_ees * num_dims * sizeof(char) );
cudaMalloc((void **) &d_measure, num_gg * sizeof(double) );
cudaMalloc((void **) &d_dot_products, num_ees * num_gg * sizeof(long) );
cudaMalloc((void **) &d_gg_x, num_ees * num_gg * sizeof(double) );
cudaMalloc((void **) &d_pp_x, num_ees * num_segments * sizeof(double) );
…
int threadsPerBlock = 512;
int blocksPerGrid =(num_gg + threadsPerBlock - 1) / threadsPerBlock;
…
GetMeasureGPU<<<blocksPerGrid, threadsPerBlock>>>(…);
…
global void GetMeasureGPU(char* d_data_gg, char* d_data_pp, double* d_gg_x, double* d_pp_x,
long* d_dot_products, double d_measure,
long num_ees, long num_dims, long ff, long num_segments, long num_gg, long num_step)
{
long g = blockDim.x * blockIdx.x + threadIdx.x;
if (g < num_gg)
{
for (long e=0; e<num_ees; e++)
{
char d_data_pp_ee = d_data_pp + e * num_dims + ff;
char* d_data_gg_ee = d_data_gg + e * num_step * num_gg + gnum_step;
double d_pp_x_ee = d_pp_x[e * num_segments + ff/num_step];
double d_gg_x_ee = d_gg_x[e * num_gg + g];
for (long j=0; j<num_step; j++)
{
d_dot_products[enum_gg + g] += (long)d_data_gg_ee[j] * (long)d_data_pp_ee[j];
}
d_measure[g] += d_dot_products[e*num_gg + g] * d_gg_x_ee * d_pp_x_ee;
}
}
}
楼主，你应该先把你的问题叙述一下，这样我们边看你的内核，再根据你需要实现的问题，才能更好的考虑如何设计你的内核以及分析你现在内核可能可以优化的问题！
代码可以重写为：
for (long e=0; e<num_ees; e++)
{
for (long j=0; j<num_step; j++)
{
d_dot_products[enum_gg + g] += (long)d_data_gg[e * num_step * num_gg + gnum_step + j] * (long)d_data_pp[e * num_dims + ff + j];
}
d_measure[g] += d_dot_products[e*num_gg + g] * d_pp_x[e * num_segments + ff/num_step] * d_gg_x[e * num_gg + g];
}
简单地说:
内层循环中的代码是：分段累进（16个元素）计算两个char数组的点积；
外层循环中的代码是：将点积的结果乘以2个double值的结果叠加到d_measure上。
其实也不算复杂，主要就是这两行代码。
田园已经在2#说的很清楚了，考虑到你顶贴不易，那么我再重复一遍：
“楼主，你应该先把你的问题叙述一下，这样我们边看你的内核，再根据你需要实现的问题，才能更好的考虑如何设计你的内核以及分析你现在内核可能可以优化的问题！”
我们这个论坛是帮助用户解决问题的。但不是帮忙给出最后成果的。
我希望的是楼主不要告诉我们你的代码是干什么的，而是你要实现的具体问题是什么，因为光看代码的话，我们也很难找出如何去优化，或许你的代码执行效率已经很好了，如果是这样的话，光从内核角度是无法再有效提高的，而是要根据你想实现的问题来考虑更合适的算法！
具体的问题我也没法说清楚，这是我的一个很复杂的算法中关键的一小部分。
算法的改进应该更难了。
如果仅仅针对该内核代码的优化，是否还有改进的空间？
比如：
- 数组数据是否可以重新排列从而得到更高的效率？
- 内核的循环是否能够更好地展开？
- 是否可以使用共享内存或者常量内存？
非常感谢！！！
如果仅针对内核优化，楼主可以先用profile跑一次你的程序，然后分析一下是否有优化空间！当然如果能引入共享内存的使用当然最好，比如从你上面的代码来看，d_dot_products这个数组会频繁写入，如果能放到共享内存，那么肯定是可以提高一定的效率！此外，在某些情况下重新排列你的数组是有可能提高你的算法效率，但是从你上面的代码来看，基本上线程都是连续存取，所以排列数组看似无任何意义了！内部的循环能否展开，还是要看你需要实现的功能是否支持（比如有某种算法能够将双重循环替换），如果不行，那么很遗憾，没有办法！当然K20上有动态并行，如果算法可以，那么可以考虑用动态并行替换内层循环！
LZ春节好，我来大致说一下我的观点，供您参考。
如您本楼和顶楼的说法，您的问题主要是循环累加如何进行并行化的问题。
一般有“原子操作”和“规约操作”两个想法，这里建议您尝试一下 规约操作。
如您本楼给出的部分代码，实际上是让一个线程为一个d_measure[g]的结果做二重循环操作。(先不说这里面有访存的效率问题)那么，如果我们使用num_ees个线程来完成同样的操作，每个线程自己做一个16次循环的访存操作（注意，这样按照你代码里面的情况，访存效率依然是有问题的，但先不细说。），那么每个线程就会得到一个d_dot_products的值，并可以以此计算一个d_measure的值，每个线程的值其实对应每次e循环的工作。然后对num_ees个线程各自得到的d_measure的值进行累加，这可以使用规约加法实现。（规约加法的实现细则详见手册和各种资料）参考您1楼的代码，您原来一个线程的活，现在使用了num_ees个线程实现了，总的线程数量也应该扩大为num_ees倍。
换言之，原先一个线程干的活，现在num_ees个线程为您服务，并行性有所增强。和原来相比，内循环没变，外循环的累加变成了规约加法。
如果您理解了上述思路，并且理解了用规约加法实现累加的原理。让我们继续。
在前面的实现中，有一个问题在于访存效率的低下。
为何会这样呢？因为您的kernel中是每个线程在循环中依次访问连续的地址空间。——这种做法在CPU实现上是合理的，高效的，但是在GPU实现上则不然。GPU的结构决定了需要编号连续的线程一起访问连续的（或者放宽说局部的）地址空间（也就是0#线程访问a[0]，1#线程访问a[1]这种）。原因在于一个warp的线程是一起进行访存的，而他们的访存行为如果能满足上述要求（即“合并访问”（大体这样，细节不完全，请参阅手册）），可以一次性完成访存，这一次性拿到的数据必然是某个访存粒度（容量）的连续地址的内容。（如您不熟悉硬件行为，可以只记住本段中关于相邻线程访存最好连续的大致观点即可）
回到您的代码中，您的写法为：“ d_dot_products[enum_gg + g] += (long)d_data_gg_ee[j] * (long)d_data_pp_ee[j];”
这里面左侧的d_dot_products[enum_gg + g]是合并的，这包括一次读取和一次写入（因为累加）。右侧的两项则不然，每个线程自己每次访存是连续的，而相邻的两个线程每次读取的位置相差很远。这将造成每次访存都要丢弃大量的无效数据，访存效率可能会有一个量级的下降。
如何解决呢？
第一也还可以考虑将内层循环用规约求和展开的方法，16个线程代替之前一个线程的工作，这样一个warp可以代替原来两个线程的工作，得到两个d_dot_products的数据，进而得到两个d_measure的数据，然后就地自己相加，然后再参加总体的规约。这样您需要num_ees/2个warp为您服务。
这样解决的好处在于搞定的访存，但是缺点在于，内循环的规约算着算着，每个warp里面就只有十分少量的线程在干活了，这一点值得改进。
第二，您也可以考虑根据您的算法实现意图修改，比如说重新划分任务，再比如说也可以考虑改变数据的安排格式什么的。关于这些我就不再进行进一步的推测了，只是建议您一个总体的思想，那就是，累加可以用并行规约的方式解决。
以上是个人的一点观点，根据代码反推实现意图经常容易出错，还请LZ和其他人多加补充和指正。
目测本帖是蛇年沙发帖，恭祝LZ（帮主？）癸巳蛇年大吉，万事如意！
system
10
另外，诚如上面TY版主反复说过的那样，您可以详细阐述一下您的算法意图，或者您自己在了解CUDA实现特点的情况下，做算法层面上的划分和改动，以期达到更好的效果。
同时，对于具体的实现，建议您跑nvvp（nvidia visual profiler）看看具体有什么瓶颈，从而有针对性地优化和改动。
祝您蛇年大吉，编码顺利，BUG退散，风轻云淡~
system
11
谢谢ice！
请问，“将内层循环用规约求和展开的方法，16个线程代替之前一个线程的工作，这样一个warp可以代替原来两个线程的工作，得到两个d_dot_products的数据，进而得到两个d_measure的数据，然后就地自己相加，然后再参加总体的规约。”能否给出代码以便于理解？
这里num_ees = 12，num_step=16，都比较小，跟一般的合并访问和归约求和的例子还是不太一样。
非常感谢！
祝春节愉快！
system
12
LZ您好，这里的实现应该和普通的规约实现差别不大，您可以参考一下规约计算例子里面最后计算thread数量少于一个warp的时候的实现代码。
根据前面9#的说法，在num_ees=12时，3#的循环可以使用num_ees/2=6个warp=192threads为您服务，192 个threads 组成一个block的话，仅从线程数量上讲，还是比较合适的。
最后根据您顶楼的代码，您可以使用num_gg个block为您服务，即之前一个thread的工作，现在一个block搞定。如果num_gg不是很小的话，这样用还是有一定意义的。
另外，依然建议您自己实现这个算法，以及，如果从算法高度全盘考虑，或许有更好的效果。
以上观点供您参考，祝您春节愉快！
system
13
ice您好！
看了几天也没搞明白，直看得非常上火、嘴巴起泡，CUDA编程实在太难懂了，找不到由浅入深的教材。
您能不能附上简单的代码，以便更好地理解！
非常感谢！！！
祝元宵节快乐！！！
system
14
LZ您好，冰冻三尺非一日之寒，如果您需要深入浅出地学习CUDA，一些国外大学的CUDA课程的视频资料是不可多得的优秀学习资源，建议您参考。此外，官方附带的手册，常见的教材以及论坛，讨论群等都可供您参考。
并行规约例子真的是最为常见的CUDA例子的，可以说到处都是，您可以在此基础上修改以适应您的算法。只有您自己弄懂，调好的代码，才对您解决问题有实际意义，甚至您可以站在算法的高度上重新设计实现。
祝您元宵节快乐！
system
15
"国外大学的CUDA课程的视频资料"哪里有啊？
你写个简单的代码给我看看，我不是更容易懂了吗？纸上谈兵总是肤浅呀。
谢谢！
system
16
"国外大学的CUDA课程的视频资料"哪里有啊？
你写个简单的代码给我看看，我不是更容易懂了吗？
谢谢！
system
17
论坛的版主是为您提供技术支持和咨询的，而不是您的工人。
如果您需要直接实现，或者如您在1#说的，直接给你某代码的优化版本，我建议您考虑就近的软件外包公司。
system
18
以及，这是我的个人看法，不排除有其他版主、会员、技术支持直接理解完您的代码并给出相应的优化实现。
在此我建议您修改1#, 将您的程序的代码放大字体，并等待看是否有其他版主、会员、NV技术支持上线后看到帖子，为您实现。
system
19
该类视频各大公开课网站/论坛都有的，iTUNEs U上面也有不过无中文字幕。请咨询度娘。
此外，我并不负责向您提供可用代码。
您可以根据上述讨论内容自己写出具体实现，或者寻找第三方写出该实现。
如对写好的实现有疑问，可以继续在论坛讨论。
你这个问题跟没回答一样。
也不知道你当这个版主有什么意义？混饭吃吗？