我是一名四年级大学生，正在从事并行计算课程项目。在选择正确的算法来展示 GPU 性能与其顺序对应的性能比较方面，我做出了一个非常错误的决定。我选择的是Bat优化算法，并行化如下： 线程0：负责计算平均值并通过更新标志来应用停止标准。 所有其他线程：更新蝙蝠种群的一代/迭代中每个蝙蝠的每个属性。 先看代码：
__device__ void startAlgo(Bat *bats, int N, unsigned long long seed){
__shared__ float global_best_fitness;
__shared__ float average_best_position_of_batSwarm;
__shared__ float global_best_position;
curandState *state = new curandState; // Should ideally be per-thread and persistent, not recreated in a loop
// Initialize random states once per thread
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < N) {
curand_init(seed, idx, 0, state);
}
if(threadIdx.x == 0){
setGlobalandAverage(bats, N, global_best_fitness, global_best_position, average_best_position_of_batSwarm);
while(!stopFlag){
printf("Average personal best position: %f\n",average_best_position_of_batSwarm);
float prev_avg = average_best_position_of_batSwarm;
CalculateFitnessAverage(bats, N, average_best_position_of_batSwarm);
ApplyStoppingCriteria(prev_avg, average_best_position_of_batSwarm);
// syncThreads(); // Synchronize after potentially modifying stopFlag or other shared variables
}
} else {
while(!stopFlag){ // Ensure this check is dynamic
for(int i = threadIdx.x; i < N; i += blockDim.x * gridDim.x){ // Distribute work more evenly
performWork(&bats[i], global_best_position, state);
printf("id: %d, v: %f, p: %f, f: %f, l: %f, pr: %f, fit: %f, pbfit: %f, pbp: %f\n",i,bats[i].velocity,bats[i].position,bats[i].frequency,bats[i].loudness,bats[i].pulse_rate,bats[i].fitness,bats[i].personal_best_fitness,bats[i].personal_best_position);
}
__syncthreads(); // Sync all threads to recheck the stopping flag
}
}
delete state; // Clean up the state
}
该函数在初始化后调用。
我已经确认每个线程确实进入了该函数，但发生的情况是只有线程 0 在循环中执行其工作，并且由于 Bat 群体中没有更新，因此程序停止。我还确认它与线程 0 首先启动或类似的事实无关。我什至尝试添加一个 spinwait 机制来检查其余线程是否至少开始执行一次工作，但没有成功。我意识到可能有很多菜鸟错误，但由于我在 Windows 上，我无法调试 cuda 程序。
如果您没有调试器，只需插入大量打印语句，这将允许您查看代码到达的位置。 打印语句的结构如下：
#define print(format, ...) do { printf("c T:%02i W:%02i B:%02i Line:%4i " format "", ThreadId(), WarpId(), blockIdx.x, __LINE__, __VA_ARGS__); } while (0)
您可能还想在其中添加
__FILE__NVidia GPU 确实不喜欢低于扭曲级别的线程分歧。为了获得最佳性能，warp 中的所有线程都需要执行相同的代码。或者至少扭曲中的单个线程需要执行，而其余线程则不执行任何操作。这比在 warp 中使用 2 个（或更多）线程执行不同的操作要高效得多。
同样在 Pascal 上，warp 中的线程以 lockstep 执行之前，因此如果一个线程执行
ifelseifelse如果执行一个扭曲中的线程发散的循环，则会得到非常低效的代码。您也可能会陷入僵局。在 Volta 及更高版本上，您不会遇到死锁，但仍然会效率低下，因为线程调度程序仅与 Pascal 略有不同，并且不会在同一周期内向同一经束中的线程发出不同的指令。
这是我重写你的伪代码的方法。
//init the RNG outside of the main loop
template <int blocksize> //make blocksize a compile time constant
__device__ void startAlgo(Bat *bats, int N, unsigned long long seed, curandState* state) {
__shared__ float global_best_fitness;
__shared__ float average_best_position_of_batSwarm;
__shared__ float global_best_position;
__shared__ volatile int stopFlag;
stopFlag = 0;
static_assert(blocksize > 32);
assert(blocksize == blockDim.x);
//assert(__isShared(state)); //perhaps store RNG in shared mem for speed?
// Initialize random states once per thread
const int idx = threadIdx.x; // + blockIdx.x * blocksize; //surely there is only one block?
const int idx2 = idx - 32; //split this block in two
const int blockDim2 = blocksize - 32;
if (idx < N) {
curand_init(seed, idx, 0, state);
}
const int warpid = threadIdx.x % 32; //warp0 does the avg, other warps do work.
if (warpid == 0) { //on K1000 only warps can execute different paths efficiently.
setGlobalandAverage(bats, N, global_best_fitness, global_best_position, average_best_position_of_batSwarm);
while (!stopFlag) {
printf("Average personal best position: %f\n",average_best_position_of_batSwarm);
float prev_avg = average_best_position_of_batSwarm;
CalculateFitnessAverage(bats, N, average_best_position_of_batSwarm);
const bool done = ApplyStoppingCriteria(prev_avg, average_best_position_of_batSwarm);
if (done) { atomicOr(&stopFlag, 1); }
} //while !done
} else {
assert(warpid > 0);
while (!stopFlag){ // Ensure this check is dynamic
for(int i = idx2; i < N; i += blockDim2) { // Distribute work more evenly
performWork(&bats[i], global_best_position, state, idx2); //to avoid having to recalc idx2
printf("id: %d, v: %f, p: %f, f: %f, l: %f, pr: %f, fit: %f, pbfit: %f, pbp: %f\n",i,bats[i].velocity,bats[i].position,bats[i].frequency,bats[i].loudness,bats[i].pulse_rate,bats[i].fitness,bats[i].personal_best_fitness,bats[i].personal_best_position);
}
//__syncthreads(); // Sync all threads to recheck the stopping flag
} //while !done
} //if !warp0
__syncthreads(); //now you can call synthreads, because all threads meet up here
delete state; // Clean up the state
}
通过将工作划分到不同的warp中，同一warp中的线程之间不再发生死锁，效率大大提高。
通过将
__syncthreads