游戏程序员的核心竞争力是什么？ 第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
相关话题
有哪些事实没有一定计算机知识的人不会相信？<<深度探索c++对象模型>>中的虚继承看着蛋疼,感觉这在实际中也没多大用,需要继续深究吗?
为什么美国程序员工作比中国程序员工作轻松、加班少？
如何理解互斥锁、条件锁、读写锁以及自旋锁？
应该像对待色情产业一样禁止网络游戏产业吗？
C#中关于List<T>和HashSet<T>应用的效率问题？
有哪些毁掉一个游戏的败笔？
长得丑的女程序员如何向男程序员讨教如何写代码？
你是因为什么而喜欢上编程的？
Python中 __init__的通俗解释是什么？
前一个讨论
下一个讨论
相关的话题
为什么中国没有厂家能够推出比肩 Switch 的游戏机？如何评价巫师系列的制作组cd projekt red？
Windows 的路径中表示文件层级为什么会用反斜杠 “”，而 UNIX 系统都用斜杠 “/”？
学编程一年，前两天到一家公司实习，今天主动提出离开，这两天是我第一次感觉到压力。我该怎么办？
游戏制作为什么这么耗钱，经常可以看到许多的独立工作室因资金不足额搁置游戏制作的？
如何不使用loop循环，创建一个长度为100的数组，并且每个元素的值等于它的下标？
config、option、setting这三者在程序世界里是什么区别？
在游戏内容为王的时代，开发者还要对应用商店点头哈腰吗？
你无意中发现过哪些图灵完全的系统？
计算机编程算是否是理科中比较偏文的科目？
游戏制作为什么这么耗钱，经常可以看到许多的独立工作室因资金不足额搁置游戏制作的？
如何评价游戏公司「心动」宣传片中传达的「做好产品，钱就是副产品」理念？在当下游戏环境中是否可行？
非计算机专业，想刷leetcode，请问在此之前需要做什么准备？
如果有两颗药丸，一颗吃了让你写代码100%不出错，另一颗吃了能让你100%发现并修改bug，选哪颗？
为什么好多编程“牛人”不喜欢用 Microsoft Visual Studio？
造轮子厉害是一种什么样的体验？
为什么国内的游戏公司吃相越来越难看了？
许多老程序员不建议新手用IDE集成开发环境，而是用编辑器＋编译器，用命令行编译，这个怎么看？
游戏行业真的不行了吗？
如何看待wegame宣布1000万流水以内的国产单机游戏0抽成？
iOS平台上开发游戏一般用什么音频格式做音乐音效？
如何看待 Python 之父 Guido 加盟微软？
连小孩子都在学编程，报考计算机专业还有意义吗？
psp2000推荐玩的游戏？
AI技术是否是制约单机游戏突破的瓶颈？
不懂就问：Linux、安卓和鸿蒙到底是什么关系？
初中未毕业的人自学得了编程吗？
是否存在那种，已经复杂到无法继续有效维护的软件？如果没有，哪些是最接近的？
如何评价网易推出的《荒野行动》PC版？
编程学到可以「抢月饼」的程度，要选哪一门语言入门比较容易，需要多长时间？