游戏程序员的核心竞争力是什么？ 第1页
1
yaobrother 网友的相关建议:
很简单，别人不会的你会。
比如，Bitonic Sort（双调排序，时间复杂度 ）是一种并行排序算法，特别适合在GPU上实现。我自己用OpenCL实现了一个并进行了调优，算法是我用纸笔推导的，分为5个kernel。在GPU上编程不同于CPU,没有调试器可用，也没有printf，数组越界电脑直接黑屏。
我就直接用C++实现了一个在CPU上运行的OpenCL kernel，实现了__local限定符 get_global_id get_local_id barrier等函数，用它实现的我的排序算法。这样就可以很方便的进行调试了，验证完成后移植在GPU上。
本人看过多个Bitonic Sort的实现， 在github搜索的基本上就是一堆玩具，没有我实现的好，Microsoft DirectX SDK中的也是玩具。DirectX-Graphics-Samples的 使用了Indirect模式，但是其numthreads是写死了的。 CUDA SDK Samples的 4个Kernel ,还算可以。
相比于其他的实现，我的这个使用#define设置不同的local size，在不同的硬件上实现更好的性能。
这就是Bitonic Sort的kernel代码（完全是自己实现的）：
#if VEC_STEP == 4 #define VEC_SEQ (uint4)(0,1,2,3) #define SCALAR_TYPE int #define VEC_TYPE int4 #define VEC_TYPE_AS as_int4 #define LOGICAL_TYPE uint4 #define LOGICAL_TYPE_AS as_uint4 #define INDEX_TYPE uint4 #define INDEX_TYPE_AS as_uint4 #define VLOADN vload4 #define VSTOREN vstore4 #define INPUT_TO_PRIVATE(A,B,IA,IB,INPUT) A.s0=INPUT[IA.s0]; A.s1=INPUT[IA.s1]; A.s2=INPUT[IA.s2]; A.s3=INPUT[IA.s3]; B.s0=INPUT[IB.s0]; B.s1=INPUT[IB.s1]; B.s2=INPUT[IB.s2]; B.s3=INPUT[IB.s3]; #define PRIVATE_TO_INPUT(A,B,IA,IB,INPUT) INPUT[IA.s0] = A.s0; INPUT[IA.s1] = A.s1; INPUT[IA.s2] = A.s2; INPUT[IA.s3] = A.s3; INPUT[IB.s0] = B.s0; INPUT[IB.s1] = B.s1; INPUT[IB.s2] = B.s2; INPUT[IB.s3] = B.s3; #else #error Need VEC_SEQ #endif #ifndef LOCAL_VEC_SIZE #error Need LOCAL_VEC_SIZE #endif #ifndef WG_SIZE #error Need WG_SIZE #endif //#define LOCAL_VEC_SIZE 4096 //设 global_size 为 2^m local_size为 2^n 那么 void BitnoicSortLocalStage1_1(__local SCALAR_TYPE* input,const unsigned int num_local) { size_t gid = get_group_id(0); size_t num_groups = get_num_groups(0); size_t lid = get_local_id(0); size_t lsize = get_local_size(0); #pragma unroll for (unsigned int step1 = 1; step1 < LOCAL_VEC_SIZE; step1 <<= 1) { unsigned int sh = 0; for (unsigned int step2 = step1; step2 >= 1; step2 >>= 1) { for (unsigned int loop = lid; loop < LOCAL_VEC_SIZE / 2; loop += lsize*VEC_STEP) { INDEX_TYPE iloop = VEC_SEQ * (uint)lsize + loop; INDEX_TYPE gloop = iloop + (LOCAL_VEC_SIZE / 2)*num_local; INDEX_TYPE t1 = gloop & step1; INDEX_TYPE t2 = t1 ^ step1; INDEX_TYPE u1 = (iloop&(~(step2 - 1))) << 1; INDEX_TYPE u2 = (step2 - 1)&iloop; INDEX_TYPE a = u1 + u2 + (t1 >> sh); INDEX_TYPE b = u1 + u2 + (t2 >> sh); VEC_TYPE ia;// = input[a]; VEC_TYPE ib;// = input[b]; INPUT_TO_PRIVATE(ia,ib,a,b,input); LOGICAL_TYPE icomp = LOGICAL_TYPE_AS(ia < ib); LOGICAL_TYPE ncomp = LOGICAL_TYPE_AS(!icomp); VEC_TYPE oa=select(ia,ib,icomp); VEC_TYPE ob=select(ia,ib,ncomp); PRIVATE_TO_INPUT(oa,ob,a,b,input); } barrier( CLK_LOCAL_MEM_FENCE); sh++; } } } void BitnoicSortLocalStage1_2(__local SCALAR_TYPE* input,const unsigned int step1_begin,const unsigned int shift,const unsigned int num_local) { size_t gid = get_group_id(0); size_t num_groups = get_num_groups(0); size_t lid = get_local_id(0); size_t lsize = get_local_size(0); unsigned int sh = shift; unsigned int step1 = step1_begin; #pragma unroll for (unsigned int step2 = (LOCAL_VEC_SIZE / 2); step2 >= 1; step2 >>= 1) { for (unsigned int loop = lid; loop < LOCAL_VEC_SIZE / 2; loop += lsize*VEC_STEP) { INDEX_TYPE iloop = VEC_SEQ * (uint)lsize + loop; INDEX_TYPE gloop = iloop + (LOCAL_VEC_SIZE / 2)*num_local; INDEX_TYPE t1 = step1 & gloop; INDEX_TYPE t2 = t1 ^ step1; INDEX_TYPE u1 = (iloop&(~(step2 - 1))) << 1; INDEX_TYPE u2 = (int)(step2 - 1)&iloop; INDEX_TYPE a = u1 + u2 + (t1 >> sh); INDEX_TYPE b = u1 + u2 + (t2 >> sh); VEC_TYPE ia;// = input[a]; VEC_TYPE ib;// = input[b]; INPUT_TO_PRIVATE(ia,ib,a,b,input); LOGICAL_TYPE icomp = LOGICAL_TYPE_AS(ia < ib); LOGICAL_TYPE ncomp = LOGICAL_TYPE_AS(!icomp); VEC_TYPE oa=select(ia,ib,icomp); VEC_TYPE ob=select(ia,ib,ncomp); PRIVATE_TO_INPUT(oa,ob,a,b,input); } barrier(CLK_LOCAL_MEM_FENCE); sh++; } } //运行1次 __kernel __attribute__((reqd_work_group_size(WG_SIZE,1,1))) void BitonicSortLocalStage1_1Kernel( __global SCALAR_TYPE* input, const unsigned int step1, const unsigned int step2, const unsigned int len ) { size_t gid = get_group_id(0); size_t num_groups = get_num_groups(0); size_t lid = get_local_id(0); size_t lsize = get_local_size(0); unsigned int loop_count = len/LOCAL_VEC_SIZE; __local SCALAR_TYPE value[LOCAL_VEC_SIZE]; for(unsigned int ng=gid;ng< loop_count ; ng+=num_groups) { async_work_group_copy(value,input+(ng*LOCAL_VEC_SIZE) , LOCAL_VEC_SIZE,(event_t)0); barrier( CLK_LOCAL_MEM_FENCE); //wait_group_events(1,&e1); BitnoicSortLocalStage1_1(value,ng); async_work_group_copy(input+(ng*LOCAL_VEC_SIZE),value , LOCAL_VEC_SIZE,(event_t)0); barrier( CLK_LOCAL_MEM_FENCE); //wait_group_events(1,&e2); } } //运行m-n次 __kernel __attribute__((reqd_work_group_size(WG_SIZE,1,1))) void BitonicSortLocalStage1_2Kernel( __global SCALAR_TYPE* input, const unsigned int step1, const unsigned int step2, const unsigned int len, const unsigned int sh ) { size_t gid = get_group_id(0); size_t num_groups = get_num_groups(0); size_t lid = get_local_id(0); size_t lsize = get_local_size(0); unsigned int loop_count = len/LOCAL_VEC_SIZE; __local SCALAR_TYPE value[LOCAL_VEC_SIZE]; for(unsigned int ng=gid;ng< loop_count ; ng+=num_groups) { async_work_group_copy(value,input+(ng*LOCAL_VEC_SIZE) , LOCAL_VEC_SIZE,(event_t)0); barrier( CLK_LOCAL_MEM_FENCE); BitnoicSortLocalStage1_2(value,step1,sh,ng); async_work_group_copy(input+(ng*LOCAL_VEC_SIZE),value , LOCAL_VEC_SIZE,(event_t)0); barrier( CLK_LOCAL_MEM_FENCE); } } void BitonicSortLocalStage2(__local SCALAR_TYPE* input) { size_t gid = get_group_id(0); size_t num_groups = get_num_groups(0); size_t lid = get_local_id(0); size_t lsize = get_local_size(0); size_t gsize = get_global_size(0); #pragma unroll for (unsigned int step1 = LOCAL_VEC_SIZE / 2; step1 >= 1; step1 >>= 1) { for (int loop = lid; loop < LOCAL_VEC_SIZE / 2; loop += lsize*VEC_STEP) { INDEX_TYPE iloop = VEC_SEQ * (uint)lsize + loop; INDEX_TYPE u1 = (~(step1 - 1) & iloop) << 1; INDEX_TYPE u2 = ((step1 - 1) & iloop); INDEX_TYPE a = u1 + u2; INDEX_TYPE b = a + step1; VEC_TYPE ia;// = input[a]; VEC_TYPE ib;// = input[b]; INPUT_TO_PRIVATE(ia,ib,a,b,input); LOGICAL_TYPE icomp = LOGICAL_TYPE_AS(ia < ib); LOGICAL_TYPE ncomp = LOGICAL_TYPE_AS(!icomp); VEC_TYPE icomp_ = VEC_TYPE_AS(icomp & 1); VEC_TYPE ncomp_ = VEC_TYPE_AS(ncomp & 1); VEC_TYPE oa = ia*icomp_ + ib*ncomp_; VEC_TYPE ob = ia*ncomp_ + ib*icomp_; PRIVATE_TO_INPUT(oa,ob,a,b,input); } barrier(CLK_LOCAL_MEM_FENCE); } } //运行1次 __kernel __attribute__((reqd_work_group_size(WG_SIZE,1,1))) void BitonicSortLocalStage2Kernel( __global SCALAR_TYPE* input, const unsigned int step1, const unsigned int step2, const unsigned int len ) { size_t gid = get_group_id(0); size_t num_groups = get_num_groups(0); size_t lid = get_local_id(0); size_t lsize = get_local_size(0); unsigned int loop_count = len/LOCAL_VEC_SIZE; __local SCALAR_TYPE value[LOCAL_VEC_SIZE]; for(unsigned int ng=gid;ng< loop_count ; ng+=num_groups) { async_work_group_copy(value,input+(ng*LOCAL_VEC_SIZE) , LOCAL_VEC_SIZE,(event_t)0); barrier( CLK_LOCAL_MEM_FENCE); BitonicSortLocalStage2(value); async_work_group_copy(input+(ng*LOCAL_VEC_SIZE),value , LOCAL_VEC_SIZE,(event_t)0); barrier( CLK_LOCAL_MEM_FENCE); } } //运行(m-n) + (m-n)(m-n-1)/2 次 __kernel __attribute__((reqd_work_group_size(WG_SIZE,1,1))) void BitonicSortGlobalStage1Kernel( __global SCALAR_TYPE* input, const unsigned int step1, const unsigned int step2, const unsigned int len, const unsigned int sh ) { size_t gid = get_global_id(0); size_t num_groups = get_num_groups(0); size_t lid = get_local_id(0); size_t lsize = get_local_size(0); size_t gsize = get_global_size(0); for (unsigned int loop = gid*VEC_STEP; loop < len / 2; loop += gsize*VEC_STEP) { unsigned int t1 = step1 & loop; unsigned int t2 = t1 ^ step1; unsigned int u1 = ((~(step2 - 1))&loop) << 1; unsigned int u2 = (step2 - 1)&loop; unsigned int a = u1 + u2 + (t1 >> sh); unsigned int b = u1 + u2 + (t2 >> sh); VEC_TYPE ia; VEC_TYPE ib; ia = VLOADN(a / VEC_STEP, input);//input[a]; ib = VLOADN(b / VEC_STEP, input);//input[b]; LOGICAL_TYPE icomp = LOGICAL_TYPE_AS(ia < ib); LOGICAL_TYPE ncomp = LOGICAL_TYPE_AS(!icomp); VEC_TYPE oa=select(ia,ib,icomp); VEC_TYPE ob=select(ia,ib,ncomp); VSTOREN(oa, a / VEC_STEP, input);//input[a]; VSTOREN(ob, b / VEC_STEP, input);//input[b]; } } //运行m-n次 __kernel __attribute__((reqd_work_group_size(WG_SIZE,1,1))) void BitonicSortGlobalStage2Kernel( __global SCALAR_TYPE* input, const unsigned int step1, const unsigned int step2, const unsigned int len ) { size_t gid = get_global_id(0); size_t num_groups = get_num_groups(0); size_t lid = get_local_id(0); size_t lsize = get_local_size(0); size_t gsize = get_global_size(0); for (unsigned int loop = gid*VEC_STEP; loop < len / 2; loop += gsize*VEC_STEP) { unsigned int u1 = (~(step1 - 1) & loop) << 1; unsigned int u2 = ((step1 - 1) & loop); unsigned int a = u1 + u2; unsigned int b = a + step1; VEC_TYPE ia; VEC_TYPE ib; ia = VLOADN(a / VEC_STEP, input);//input[a]; ib = VLOADN(b / VEC_STEP, input);//input[b]; LOGICAL_TYPE icomp = LOGICAL_TYPE_AS(ia < ib); LOGICAL_TYPE ncomp = LOGICAL_TYPE_AS(!icomp); VEC_TYPE icomp_ = VEC_TYPE_AS(icomp & 1); VEC_TYPE ncomp_ = VEC_TYPE_AS(ncomp & 1); VEC_TYPE oa = ia*icomp_ + ib*ncomp_; VEC_TYPE ob = ia*ncomp_ + ib*icomp_; VSTOREN(oa, a / VEC_STEP, input);//input[a]; VSTOREN(ob, b / VEC_STEP, input);//input[b]; } } 后来找到工作后，将其改成DirectX Compute Shader实现的。完全独自一人完成，公司除本人外没有一个人会。这个是为了做粒子效果用的，而粒子效果需要进行排序。
以后打算使用HLSL Shader Model 6 的新指令（vote, ballot, shuffle）进行性能优化，性能可以更快。
在GPU上还有一种排序 Radix Sort（基数排序），Bullet物理引擎用的就是Radix Sort。与Bitonic Sort相比 速度会快一些，但是会占额外的内存空间（Bitonic Sort不需要，可以就地进行排序，Radix Sort需要2个Buffer，一个读一个写），浮点数比较需要处理（Bullet引擎的FlipFloat）。本人也用OpenCL实现了一个并进行了测试，包含7个Kernel。其中有ExclusiveScan,Histogram等并行计算常用算法。
=============
2019年1月20日更新：
二叉树，估计大家都知道是什么。
但在GPU上建造二叉树各位听说过吗？
Bullet物理引擎就用到了，用于射线相交检测的BVH，需要将AABB包围盒转换成Morton Code后排序，然后根据排序后的Morton Code建造二叉树。
具体原理请看
和论文 Maximizing Parallelism in the Construction of BVHs, Octrees,and k-d Trees （该链接里就有）
借用论文中的一张图：
如图
红色代表中间节点 绿色代表叶子节点。一共有8个叶子节点，7个中间节点。
水平的箭头代表二进制公共前缀的高度，如3-4之间就为0，0-1之间为3
让0为最高的节点，开始查找最高的公共前缀高度，找到3-4之间最高，于是0节点指向3，4节点
3节点 2-3公共前缀高度为4，3-4公共前缀高度为0，2-3比3-4低所以向左查找。查找比3-4高度低的最高的节点。
找到1-2之间公共前缀高度为2，最高。3节点指向1，2节点。
4节点 3-4公共前缀高度为0，4-5公共前缀高度为1，3-4比4-5高所以向右查找。查找比3-4高度低的最高的节点。
找到4-5之间公共前缀高度为1，最高。4节点指向4，5节点。由于4节点指向自身，因此指向叶子节点。
叶子节点间，中间节点间完全无依赖，可以完全并行的建造。
Bullet物理引擎使用的方法与上述有些不同。
Bullet物理引擎的做法是：
每个叶子节点，如果左边的公共前缀高度比右边大，那么就指向左边的中间节点，如果右边大就指向右边的。最左的节点 公共前缀高度最小，最右的节点公共前缀高度最大。
每个中间节点，分别从左右方向查找离该节点最近的公共前缀高度比该节点高的节点，如果找不到， 公共前缀高度为最小。
如果找到的左边的公共前缀高度比右边大，那么就指向左边的中间节点，如果右边大就指向右边的。
当时为了搞懂算法原理，我把运行时产生的数据输出成文本，才找出了其中的规律。
这就是Bullet物理引擎，一个自带的Demo，看看左下角有多少个立方体。
这个Demo中的拾取算法（射线-三角形相交）就用到了BVH，每一帧的拾取都会建造BVH。碰撞检测也可以用BVH，但这个Demo用的是SAP。
写本文时入职才2个月，第一份工作。
可以看看我的另一个回答：
//=====================
2019年1月27日更新：
看到评论中有些人用过CUDA，我就说一说CUDA 和DirectX12 做计算的区别，难度不是一个等级的：
CUDA分配Buffer 直接用 cudaMalloc 。DirectX12用CreateCommittedResource，再根据资源的使用方式创建view。
CUDA上传用cudaMemcpy，而DirectX12 有2种方法：第一种直接创建一个 upload buffer ,然后memcpy。第二种创建一个 default buffer,再创建一个upload buffer，先memcpy到upload buffer,再根据资源类型 CopyBufferRegion CopyTextureRegion 到default buffer。
Root Signature ：shader的资源绑定布局，相当于函数参数的定义，这个CUDA没有。
Descriptor Heap： CUDA没有，存放 view的地方。相当于预先定义的一个数组，存放Buffer 的指针。
Resource Barrier：资源状态转换，CUDA没有。上传资源时upload buffer状态转换为 COPY_SOURCE ，default buffer状态转换为COPY_DEST。当参数使用时default buffer状态转换为SHADER_RESOURCE。
Command Queue Command List Command Allocator ：所有的渲染命令如 Draw Dispatch Copy Clear 都要记录到 Command List中，而 Command List需要指定Command Allocator 。完成后Close，Command Queue 再ExecuteCommandList执行记录的命令。
Fence Signal ：ExecuteCommandList 是非阻塞的，执行后立即返回。需要等待的话创建一个Fence, 然后Command Queue再Signal创建的Fence，Fence再Wait。
yu-gu-47 网友的相关建议:
11年跑步，17年开始用软件记录。跑步给我最大的改变是对路程的认知。5公里的路程第一反应是跑过去，10公里的路程，就想骑自行车也就40分钟的事，而不是第一反应是坐公交或者打车。还有，去到一个陌生的城市，我也会用跑步打开对这座城市的了解。那种走路太慢，跑步正好的惬意，感觉真的很好。最后一点，听着音乐跑步成了我的一个自我奖励点和情绪宣泄点，有啥事达到目标了，奖励自己今天听着动感的音乐跑个5公里，有郁闷的事情了，绕着操场咔咔跑10公里下来，坏情绪基本就掉没了。不过也有缺点，倘若天天跑5公里，1-2个月下来，膝盖真受不了，再就是换衣服+热身+跑步+洗澡+吃饭，这一套下来，2小时就没了，如果有实验啥的，还挺耽误时间的，以上。
1
相关话题
你在写程序的时候经历过哪些有趣的事？GitHub 或者其他的开源平台中是否有一些适合初学者的 C++ 项目？
今年25了，之前一直做cocos棋牌，我该不该降薪去转ue4？
突然想开一家程序员主题的餐馆，名字就叫程序员的菜，菜名就叫各种语言中的关键字，各位指点一哈，有前途没？
如何看待“Python星人”这个群体？
大一学生，刚开始学习编程，但感觉自己并没有天赋怎么办？
为什么要学习别人的设计模式？
如何理解 “c++缺少对象级别的消息发送机制” 这句话？
人工智能趋势下，孩子学习编程有哪些意义，对未来发展有好处吗？
为什么我们需要了解编程的历史？
前一个讨论
下一个讨论
相关的话题
如果代码按行付费，会有什么好玩的事情发生？如何系统地自学 Python？
GitHub 上可供新手阅读和玩耍的 Java 项目有哪些？
如何抽象评判现有语言优劣，继而设计一款别具优雅的计算机语言 X ？
如何看待螃蟹科技燕飞宏的声明？
Rust 语言现在什么情况，为什么知乎上不怎么讨论Rust语言了呢？
你看过/写过哪些有意思的代码？
把BAT的机房炸掉，公司是不是就垮了?
哪个瞬间你觉得中国的游戏产业真叼？
我学编程为什么难？是思维方式不对还是学习方式不对？
做游戏不仅要编程能力还要计算机其它方面的知识吗？
你见过哪些神一样的优化操作？
Unreal Engine 5（虚幻引擎5）会对游戏、影视等产业带来什么样的变革?
过度依赖框架有什么不好？
为什么有一些程序员认为开源有助于发现软件中的缺陷？
传统的try-catch异常处理是否是编程语言发展中的弯路？
编写基于机器学习的程序，有哪些编写和调试的经验和窍门？
游戏为中国创造那么多GDP，带动了大批副产业（如直播，周边）发展，为什么国家还是不肯放过游戏？
python如何将变量名转化为同名字符串？
做游戏不仅要编程能力还要计算机其它方面的知识吗？
编程中如何给变量命名？有哪些规范的做法值得学习借鉴？
有哪个高手可以解读“世界黑客编程大赛第一名的作品（97年Mekka ’97 4K Intro）”?
windows10 界面是用什么语言编写的?.NET在最新的操作系统中有哪些应用
可不可以用文言文再创造一种编程语言？
随着各个领域越来越多的人都开始学习编程，计算机专业的人优势在哪里?
C 语言如何判断等差数列？
编程那么难，为什么不弄一个大众一学就会的计算机语言呢?
2020上半年，国产独立游戏基本全军覆灭，国产独立游戏未来还有机会吗？
如何用一句话说明什么是面向对象思想？
当程序员需要天赋吗？