标签：index WebGPU 冒泡排序 secondIndex sharedData var 排序 firstIndex 着色器

大家好~本文使用WebGPU的计算着色器，实现了奇偶排序。
奇偶排序是冒泡排序的并行版本，在1996年由J Kornerup提出。它解除了每轮冒泡间的串行依赖以及每轮冒泡内部的串行依赖，使得冒泡操作可以并行执行

 

目录

• 介绍奇偶排序算法
• 分析时间复杂度
• 需求
• 初步设计
• 代码实现
• 发现问题
• 改进设计
• 相关代码实现
• 改进设计
• 相关代码实现
• 限制
• 总结
• 参考资料

 

最终版本的代码在这里

介绍奇偶排序算法

假设待排序的数组为Arr1
在奇数步中，Arr1中奇数项与相邻的右边一项比较和交换；
在偶数步中，Arr1中奇数项与相邻的左边一项比较和交换;
直到一步中没有交换项，则停止

举例来说的话，如下图所示：

在每步中，红框内的两项进行比较和交换；
直到一步中没有交换项，则停止

分析时间复杂度

与冒泡排序一样，总的比较次数不变，依然为O(n^2)次
但因为为并行执行，所以时间复杂度降低为O(log2(n^2))=O(n)

需求

排序的需求如下所示:

• 对一个包含128个数字的数组进行升序的排序

初步设计

因为数组可以两两分为64个组，每个组并行执行操作，所以计算着色器只使用一个workgroup，包含64个局部单位，每个局部单位对应一个组；

在每个局部单位中：
启动一个while循环，执行每步操作，然后同步，最后判断所有局部单位在该步骤中是否有交换操作，如果都没有的话则停止循环

“执行每步操作”时判断该步骤是奇数还是偶数步，从而取对应的两项来比较和交换

代码实现

经过上面的设计后，现在我们来实现代码
计算着色器代码如下所示：

	//64个局部单位
	const workgroupSize = 64;
	// 局部单位之间的共享变量，用于存放128个数字
	var<workgroup> sharedData: array<f32,128>;
	// 局部单位之间的共享变量，用于标志所有局部单位在该步骤中是否有交换操作（只要有任意一个局部单位在该步骤中有交换操作，则该标志为true）
	var<workgroup> isSwap: bool;
	// 局部单位之间的共享变量，用于记录步骤数，从而判断是奇数还是偶数步
	var<workgroup> stepCount: u32;
	struct BeforeSortData {
	data : array<f32, 128>
	}
	struct AfterSortData {
	data : array<f32, 128>
	}
	//待排序的数组
	@binding(0) @group(0) var<storage, read> beforeSortData : BeforeSortData;
	//排序后的数组
	@binding(1) @group(0) var<storage, read_write> afterSortData : AfterSortData;
	@compute @workgroup_size(workgroupSize, 1, 1)
	fn main(
	@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>,
	) {
	//将待排序的数据读取到共享变量中
	var index = GlobalInvocationID.x * 2;
	sharedData[index] = beforeSortData.data[index];
	sharedData[index+ 1 ] = beforeSortData.data[index + 1];
	//初始化共享变量
	isSwap = false;
	stepCount = 0;
	//同步
	workgroupBarrier();
	//开始循环
	while(true){
	var firstIndex:u32;
	var secondIndex:u32;
	//判断该步骤是奇数还是偶数步，从而得到对应的两项的序号
	//偶数步
	if(stepCount % 2 == 0){
	firstIndex = index + 1;
	secondIndex = index + 2;
	}
	//奇数步
	else{
	firstIndex = index;
	secondIndex = index + 1;
	}
	//确保没超过边界
	if(secondIndex < 128){
	//将大的一项交换到后面，从而实现升序
	if(sharedData[firstIndex] > sharedData[secondIndex]){
	var temp = sharedData[firstIndex];
	sharedData[firstIndex] = sharedData[secondIndex];
	sharedData[secondIndex] = temp;
	isSwap = true;
	}
	}
	stepCount += 1;
	workgroupBarrier();
	//如果该步骤中没有交换操作的话则停止循环
	if(!isSwap){
	break;
	}
	}
	//将排序后的数据传给返回给CPU端的Storage Buffer中，从而可在CPU端得到排序后的结果
	afterSortData.data[index] = sharedData[index];
	afterSortData.data[index + 1] = sharedData[index + 1];
	}

本来我是想像设置workgroupSize一样，将128设为const的，如下所示：

	const workgroupSize = 64;
	const itemCount = 128;
	var<workgroup> sharedData: array<f32,itemCount>;

但是运行时会报错！照理说根据WGSL的文档，是不应该报错的！所以不清楚是我没搞清楚还是WGSL目前的bug？
（另外，如果使用override而不是const，也会报错！）

如果改为使用workgroupSize，则不会报错。代码如下所示：

	const workgroupSize = 64;
	var<workgroup> sharedData: array<f32,workgroupSize>;

这是因为workgroupSize在@workgroup_size中使用了。代码如下所示：

@compute @workgroup_size(workgroupSize, 1, 1)

发现问题

运行代码后，会报警告：

	1 warning(s) generated while compiling the shader:
	:74:5 warning: 'workgroupBarrier' must only be called from uniform control flow
	workgroupBarrier();
	^^^^^^^^^^^^^^^^
	:77:5 note: control flow depends on non-uniform value
	if(!isSwap){
	^^
	:77:9 note: reading from workgroup storage variable 'isSwap' may result in a non-uniform value
	if(!isSwap){
	^^^^^^

这是因为WGSL会进行Uniformity analysis检查，确保像“workgroupBarrier”这种barries是在uniform control flow中安全地调用
WGSL在检查时发现：因为isSwap被多个局部单位读写，所以为"non-uniform value"，导致所在的control flow为non-uniform

更多关于Uniformity的资料在这里：
uniformity
Add the uniformity analysis to the WGSL spec
uniformity issues

改进设计

现在需要去掉isSwap的if判断
因为isSwap的if判断是用来结束循环的，那么在去掉它之后我们就需要新的结束条件
因为总共有128个数字要排序，所以最多进行128步即可完成所有的排序

相关代码实现

所以去掉isSwap，把循环终止条件修改下，并且重构一下代码
相关代码改为：

	fn _swap(firstIndex:u32, secondIndex:u32){
	var temp = sharedData[firstIndex];
	sharedData[firstIndex] = sharedData[secondIndex];
	sharedData[secondIndex] = temp;
	}
	fn _oddSort(index:u32) {
	var firstIndex = index;
	var secondIndex = index + 1;
	if(sharedData[firstIndex] > sharedData[secondIndex]){
	_swap(firstIndex, secondIndex);
	}
	}
	fn _evenSort(index:u32) {
	var firstIndex = index + 1;
	var secondIndex = index + 2;
	if(secondIndex <128 && sharedData[firstIndex] > sharedData[secondIndex]){
	_swap(firstIndex, secondIndex);
	}
	}
	@compute @workgroup_size(workgroupSize, 1, 1)
	fn main(
	@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>,
	) {
	...
	var firstIndex:u32;
	var secondIndex:u32;
	for (var i: u32 = 0; i < 128; i += 1) {
	//偶数步
	if(stepCount % 2 == 0){
	_evenSort(index);
	}
	//奇数步
	else{
	_oddSort(index);
	}
	stepCount +=1 ;
	workgroupBarrier();
	}
	...
	}

现在就能够正确运行了

改进设计

现在我们通过判断步骤数是偶数还是奇数来进行对应的排序，这会造成“Warp Divergence”的优化问题：
不同的局部单位会进入不同的分支（偶数步或者奇数步），造成同一时刻除了正在执行的分支以外，其余分支都被阻塞了，十分影响性能
如下图所示：

参考资料：
CUDA编程——Warp Divergence

所以我们可以去掉stepCount和相关的判断，改为在每次循环中分别执行奇数排序和偶数排序，并把循环次数减半

相关代码实现

修改后的相关代码为：

	for (var i: u32 = 0; i < 64; i += 1) {
	_oddSort(index);
	workgroupBarrier();
	_evenSort(index);
	workgroupBarrier();
	}

限制

现在程序有如下的限制：

• 只有一个workgroup，意味着只能排序最多“局部单位数量*2”个数字

总结

感谢大家的学习~

计算着色器是SIMT架构，也就是说每条指令都是并行执行的。这与CPU的串行思维不同，所以我们要切换为并行的思维，并在需要同步的时候同步

另外，计算着色器的代码因为处在并行环境下，所以需要仔细优化
有下面的一些优化建议：

• 减少Warp Divergence
  • 尽量减少if判断
  • if中尽量使用常数判断，如if(const value < 2)
• 减少Bank Conflict
  • 对共享变量内存尽量连续地读写
    如对本文的共享数组sharedData，尽量连续的读写相邻的序号
• 减少同步操作
  • 如果已知确定的循环次数，可以展开循环，这样可以减少同步操作和循环的指令开销

参考资料

啥是Parallel Reduction
CUDA(六). 从并行排序方法理解并行化思维——冒泡、归并、双调排序的GPU实现
uniformity
CUDA编程——Warp Divergence

 

欢迎来到Wonder~

扫码加入我的QQ群：

扫码加入免费知识星球-YYC的Web3D旅程：

扫码关注微信公众号：

  转 https://www.cnblogs.com/chaogex/p/16578547.html

标签：index,WebGPU,冒泡排序,secondIndex,sharedData,var,排序,firstIndex,着色器
来源： https://www.cnblogs.com/wl-blog/p/16591845.html

本站声明： 1. iCode9 技术分享网（下文简称本站）提供的所有内容，仅供技术学习、探讨和分享；
2. 关于本站的所有留言、评论、转载及引用，纯属内容发起人的个人观点，与本站观点和立场无关；
3. 关于本站的所有言论和文字，纯属内容发起人的个人观点，与本站观点和立场无关；
4. 本站文章均是网友提供，不完全保证技术分享内容的完整性、准确性、时效性、风险性和版权归属；如您发现该文章侵犯了您的权益，可联系我们第一时间进行删除；
5. 本站为非盈利性的个人网站，所有内容不会用来进行牟利，也不会利用任何形式的广告来间接获益，纯粹是为了广大技术爱好者提供技术内容和技术思想的分享性交流网站。