写在前面
说是基础学习，其实只是对自己的大作业做一个简单的问题汇总，把源码放上来以供自己后续学习回顾。
先放源码：https://github.com/wen1127420395/CUDA_learning
1.任务概述
1.1任务介绍
利用并行遗传算法GA计算求解Flexible-Job-Shop-Problem。可以组队完成，每队最多3人，可以1人自己独立完成。要求用CUDA编写。提前熟悉一下GA算法和FJSP问题。
1.2任务目标
根据CUDA程序的性能和加速比得出期末成绩。
1.3语言
只能用CUDA C/C++实现
1.4教师提供参考资料（我没看）
个人网盘
1.5验收方式
1）.提交的文档报告需要有封面，封面内容包括课程名字、小组的成员与学号、提交的日期
2）.实验数据文档中必须有5组的对比数据，请自行选择最近发的python工程里面的数据
3）.在周四前大家务必将代码也考到服务器上能编译运行，验收的时候直接在这台服务器上跑。服务器是 ssh ****@127.0.0.1 密码：**** （各组协商使用，不要同时使用）
验收的时候主要看运行的时间和解的质量，也会对代码的完整性、鲁棒性进行考虑
2.调试过程
开发过程中共编写了四个版本的FJSP_CU，分别进行介绍。
其中CPU版本在1000次迭代中共用时17s。
2.1 FJSP_CU:
2.1.1思想及性能分析：
此版本仅对计算每一代所有个体性状时，即计算每个个体的最短工序时间，进行了并行计算
该版本在1000次迭代中共耗时35s左右。相比CPU版本反而差了很多。
2.1.2核心代码：
__global__ void Calculate(ob **a, good *s, int *facx, int *facy, int n) {
int x = 0;
if (n == 101) {
x = blockDim.x * blockIdx.x + threadIdx.x + 1;
}
else
x = blockDim.x * blockIdx.x + threadIdx.x + 101;
int sx[109];
int sy[109];
int mt[109];
int gt[109];
int tot = 50;
int Max = 0; //Max最大时间
if (x < n) {
for (int i = 1; i <= tot; i++) //该循环就做下面两件事
{
sx[s[x].count[i]] = facx[i]; //第几个工序对应属于第几个物件
sy[s[x].count[i]] = facy[i]; //这个工序对应物件的第几个工序
}
for (int i = 1; i <= tot; i++)
{
int tmp = 0;
int chosen;
int time = 100000;
int len = sizeof(a[sx[i]][sy[i]].m) / sizeof(int);
for (int j = 0; j < len; j++)
{
if (mt[a[sx[i]][sy[i]].m[j]] < time)
{
time = mt[a[sx[i]][sy[i]].m[j]];
chosen = a[sx[i]][sy[i]].m[j];
}
}
tmp = max(mt[chosen], gt[sx[i]]);
tmp += a[sx[i]][sy[i]].t;
Max = max(tmp, Max);
mt[chosen] = tmp;
gt[sx[i]] = tmp;
}
s[x].ans = Max;
printf("%d\t", Max);
}
}
2.1.3问题汇总：
1.sx[s[x].count[i]] = facx[i];中无法正常执行，导致最后调试总是发现s[x].ans = Max不能执行成功并返回给CPU
A:CUDA中是可以使用printf()函数的，因此调试虽然较慢但是还是可以找到问题所在的。最终发现，由于s[x].count[i]中部分无法访问到内存，则在sx[]索引的时候出错，导致cudaError异常，主要出错原因还是写串行版本的FJSP没有把内存问题处理好，到了GPU里自然出问题。
2.整个程序执行时间很长，相比CPU版本还慢了一倍还多
A：原因有二，第一，也是最关键的原因，是助教老师提醒我的，在kernel中不要使用malloc()，会非常慢，具体原因我还没有分析，而我恰恰在一开始使用了malloc()，并free。第二，由于只是对计算最短时间进行了并行计算，首先这一部分在串行版本中占比也不算很大，就算完全并行这一部分提升也不会太大，其次，由于每一代的遗传中，后一代的交叉变异等都需要前一代的数据，因此计算后必须把计算好的数据再拷贝回CPU，每次计算前又要把数据拷贝给GPU，这样在每一次迭代中都有两次大数据的CPU、GPU来回拷贝，浪费了大量的时间。
3.计算的s.ans（即最后的最短时间）有误，随代数增加不断上涨
A:原因在于没有清空每次声明的四个数组，使用memset清空mt、gt、sx、sy数组即可
2.2FJSP_CU2
2.2.1算法思想及性能分析
考虑到上一个版本在拷贝中花费了大量的时间，我们思考把数据全部存放在GPU内存中，整个遗传算法也放到GPU的一个线程中，需要用到数据时直接用传递的指针读数据即可，这样在需要并行计算时再调用其他的kernel（动态并行），这样就消灭了来回拷贝数据的时间。
想法似乎没问题，但是还是那句话，理论应该先行，不然实践结果会很难看，而且白费功夫。
最后结果是共花费了三百多秒，相比CPU版本慢了整整20~30倍，这是不能忍受的，事实上也是自己没有深入学习GPU设备的架构，不了解所谓的动态并行，在助教的帮助下发现这样的调用会非常慢，因此需要全部翻篇，重新构想算法流程。
2.2.2核心代码
这里我就不贴代码了，因为基本都是一些废代码。
2.2.3问题汇总
问题很明显，这么慢的GPU计算是不能忍受的，只能换一种想法，而之所以这么慢原因也介绍过了，总之，不要--global--中调用其他的--global--，会非常慢，只能调用其他的--device--。
2.3FJSP_CU3
2.3.1算法思想及性能分析
前面的成果全部打翻，从零开始构想新的方案。虽然代码不能使用，但是心里的算法思想已经有了，前面的一些实践内容也不会完全白费。
之所以慢，慢在了来回拷贝的问题上，那么除了把遗传算法完全塞进GPU还能用什么方法能避免来回拷贝？
其实答案也很简单，CPU仍然执行每一代的迭代过程，不同的是，在每一代中，调用并行的交叉、并行的变异、时间计算、选择等等，数据完全交给GPU，不拷贝回CPU数据。即尽量减少GPU和CPU的数据交互，一开始把GPU需要的数据全部送给他，在CPU中调用各个不同的kernel即可。同时尽可能的增加并行程度。
性能非常卓越，依然是我们自己的最初的测试数据，1000代跑到了4s左右的成绩，但代码不完整，输出错误，有些功能没有实现，例如qsort()等。
2.3.2核心代码
while (counter <= Lim)
{
printf("counter = %d\n", counter);
//变异
critical = Cross + 1;
variation << < 1, 512 >> > (s, tot, cmper, localState, critical);
cudaDeviceSynchronize();
//交叉
critical = sizeT + Cross + 1;
cross << < 1, 512 >> > (s, delta, cmper, tot, localState, critical);
cudaDeviceSynchronize();
critical = sizeT + Cross * 2 + 1;
//critical = 101;
// 执行kernel
Calculate << < 2, 512 >> > (a, s, facx, facy, critical);
cudaDeviceSynchronize();
//遗传
genetic_0 << < 2, 512 >> > (best, critical);
//qsort(best + 1, (sizeT + Cross) * 2, sizeof(int), cmp2);//?
genetic_1 << < 1, 128 >> > (tmp, s, best, critical);
genetic_2 << < 1, 128 >> > (s, tmp, tot, critical);
cudaDeviceSynchronize();
counter++;
}
error = cudaGetLastError();
printf("CUDA error: %s\n", cudaGetErrorString(error));
2.3.3问题汇总
1.GPU中怎么实现qsort()函数
A：因为GPU中没有现成的qsort()重载函数使用，只能自己实现功能，一开始我打算使用别人自己实现的qsort()代码，后来发现代码理解起来比较困难，而且是递归的，非常难以调试，因此我自己写了几个冒泡排序，效率不是很高，但代码简洁易懂，非常易于调试。
这里贴上两个，第一个是用来按cmp的base偏移值给base排序，第二个是简单直接排序。
__device__ void bubbleSort_1(int *base, int tot, double *cmp) {
int tmp = 0;
for (int i = 0; i < tot - 1; i++) {
for (int j = i; j < tot - 1; j++) {
if (cmp[base[j]] > cmp[base[j + 1]]) {
tmp = base[j];
base[j] = base[j + 1];
base[j + 1] = tmp;
}
}
}
}
__device__ void bubbleSort_2(int *base, int tot) {
int tmp = 0;
for (int i = 0; i < tot - 1; i++) {
for (int j = i; j < tot - 1; j++) {
if (base[j] > base[j + 1]) {
tmp = base[j];
base[j] = base[j + 1];
base[j + 1] = tmp;
}
}
}
}
2.遗传的选择过程非常慢
A：问题出在了选择过程时需要对一个种群交叉变异后的全部个体进行性状的排序，选择排名前sizeT个进行下一代遗传，结果发现如果把这些全塞到GPU中，由于要给一个数量上千的数据进行冒泡排序，而且只能串行，会慢到难以想象，因此只好采取折中的办法，把数据拷贝回来给CPU进行排序，结果发现这样并不是很慢，出乎了我的意料。
贴一下代码：
for (int i = 1; i <= (sizeT + Cross) * 2; i++) {
best[i] = i;
}
cudaMemcpy(s, d_s, 2009 * sizeof(good), cudaMemcpyDeviceToHost);
qsort(best + 1, (sizeT + Cross) * 2, sizeof(int), cmp2);
cudaMemcpy(d_best, best, 1509 * sizeof(int), cudaMemcpyHostToDevice);
critical = sizeT + Cross + 1;
/*for (int i = 1; i <= 1000; i++) {
printf("%d\t", s[i].ans);
}*/
genetic_1 << < 1, 128 >> > (d_tmp, d_s, d_best, critical);
cudaDeviceSynchronize();
genetic_2 << < 1, 128 >> > (d_s, d_tmp, d_tot, critical);
cudaDeviceSynchronize();
3.功能全都实现了，但是输出结果不正确，具体体现：最后输出的排序前sizeT名最短时间都是0，且只有前sizeT个如此。
A：问题也很快就找到了，在calculate计算最短时间中，线程号的处理上。由于最初初始化种群时和每一代遗传过程中调用calculate需要开的线程数不一样多，因此需要对不同情况的线程号x进行不同的处理，同时也要处理临界值critical。代码如下：
__global__ void Calculate(ob **a, good *s, int *facx, int *facy, int *tot, int n) {
int x = 0;
if (n == 101) {
x = blockDim.x * blockIdx.x + threadIdx.x + 1;
}
else {
x = blockDim.x * blockIdx.x + threadIdx.x + 101;
n += 100;
}
int Max = 0; //Max最大时间
if (x < n) {
//省略了处理的内容
}
2.4FJSP_CU4
由于上一个版本已经把问题都解决了，输出也正确，问题在于遗传算法似乎有些问题，会陷入局部最优，而且还要改一下读取老师给的测试数据的部分代码。最终把代码进行优化，让代码格式更规范一下。
