第二届CUDA On Arm夏令营考题详解：挂了一半的考题究竟长啥样？

啥也不说了，上题！

矩阵操作

要求，输入一个1000x1000的方阵：判定如果x,y坐标都为偶数的元素，则对该元素做平方，否则该元素值减一。

这个题目需要申请输入矩阵元素个数的线程，每个线程读取输入矩阵的一个元素，并判定其坐标，对该元素进行处理，最后将结果写入输出矩阵

解题：

这道题比较简单，大部分同学都过了。

InputMatrix

OutputMatrix

Grid

求向量Top2问题

实现求包含1000000个元素的数组的最大的两个值

解题：

1.注意我们在这里的整体思想是：

a.每个线程只干两件事，把这个线程读入的数据中最大的两个值按照顺序写入shared mem

b.与其他线程合作，求出它所在的block内shared mem中所有数据的两个最大值，并写入输出向量

2.如下图所示，按照流程，我们先将1M个数据传给GPU，每个线程先读取一个数据，并将数据写入shared mem。

3.这里的top1和top2分别为最大的值和第二大的值4.在main函数中第一次调用kernel func时，每个线程中的top1为它读取的数，top2为INT_MIN

4.接下来就是一个Block内部的事情

5.此时，在图中的例子中，每个block中的sharedmem中有16个数据，每个block有8个线程

6.调用一半的线程，每个线程读取四个数据，从中选出两个最大的值4.把这两个最大的值写入sharedmem中该线程对应的位置，

所以每个线程读取的四个数据为：这里的lich指的是sharedmem

lich[2 * threadIdx.x]

lich[2 * threadIdx.x + 1]

lich[2 *(threadIdx.x + i)]

lich[2 *(threadIdx.x + i) +1]

在这一半的线程完成任务后，同步一下，不要出现读写冲突问题

完成之后，每个Block中sharedmem中数据已更新。

上图为blockIdx.x==0的线程块的例子，其他的线程块也是以相同的方式处理

注意：此时红色框中部分代码是用来选取四个值中的最大的两个值。这里我们在上一步和下一步都会将较大的数放在前面的位置，所以已知：

lich[2 * threadIdx.x]> lich[2 * threadIdx.x + 1]

lich[2 *(threadIdx.x + i)]> lich[2 *(threadIdx.x + i) +1]

7.接下来我们继续迭代，使用上一步线程数一半的线程，继续处理剩下的数据2.这里红色框中的部分是有点绕的，基本逻辑是：

已知：a>b,c>d

则：top1= max(a,c)//这个没有争议

top2 = min(max(a,d),max(b,c))

top2如果还不明白，可以展开为if的方式：

if(a<d)return min(d,c);//top2==d

elseif (b>c) return min(a,b);//top2 == b

elseif (b<c) return min(a,c);//top2 == min(a,c)

此时，每个Block中的sharedmem中数据又更新了

8.接下来我们继续迭代，使用上一步线程数一半的线程，继续处理剩下的数据

9.此时，我们进行该for循环中最后一个迭代

当这一步更新后，每个线程块的sharedmem中前两个位置保存的就是这个线程块处理的所有数据中最大的两个值

10.接下来，每个线程块的threadIdx.x==0的线程将它所在block中sharedmem前两个位置的数据，即：top1和 top2，读取出来，按照线程块的序号，写入输出数组中

11. 接下来，我们第二次调用核函数，重复之前的操作。

注意此时，我们只调用了一个线程块，如果线程不够用，会触发红色框中代码，每次加一个stride，直到处理完所有数据。每个线程还是将它所读到最大的两个值，放在对应的sharedmem中

矩阵转置

按照矩阵转置的公式，我们设定（按照下图所示）

输入矩阵为: A[16][16] M=16

输出矩阵为: A’[16][16]

保证：A ‘[y][x] = A[x][y]

为了保证在合并访存章节中讲到的，连续的线程读写连续的数据

我们利用shared mem进行优化

而题目要求在原本的输入矩阵空间中操作，那么我们就可以

不调用所有的线程就可以完成。

如图中所示，即调用黄色，蓝色和红色所覆盖的线程块即可

下图中所示格子中数字坐标为输入矩阵中数据值在整个矩阵中的坐标位置。

不同颜色代表不同的线程block

因页面空间有限，后面就不在赘述

解题：

整体的思路很简单，下图中黄色线为一个对角线，所有黄色线下面的线程块，即blockIdx.y> blockIdx.x的线程块，处理自己本身的数据和它中心点对称的数据。即蓝色覆盖的线程块，处理他自己和绿色覆盖数据

我们以下面红色框中的线程为例，他的参数为：

blockIdx.x== 0;

blockIdx.y== 1;

threadIdx.x== 3;

threadIdx.y== 1;

注意此时绿色框内的数据皆为红色框所示线程数据

1.所以红色框中的线程，要先将(9,3)号数据和(1,11)中的数据读取出来，分别放到tile_s和tile_d中

2.再将(9,3)和(1,11)中所需要的数据写入

3.注意这里一定要判定是否越界

4.读完数据之后要记得同步，不要出现写后读产生的冲突

而黄色线穿过的线程块，即黄色和红色所覆盖的线程块，就比较简单了，只需要在块内将数据调换即可。

下图中所示格子中数字坐标为输入矩阵中数据值在整个矩阵中的坐标位置。

不同颜色代表不同的线程block

好了，以上是三道题目的解答，不知道你觉得简单？还是难？