StanFord CS149 Assignment3 简记

Github

Part 1

实现简单, 见代码

问题 1

CUDA

Running 3 timing tests:
Effective BW by CUDA saxpy:  125.259 ms         [8.922 GB/s]
saxpy without data transfer: 6.136 ms
Effective BW by CUDA saxpy:  126.499 ms         [8.835 GB/s]
saxpy without data transfer: 6.198 ms
Effective BW by CUDA saxpy:  119.858 ms         [9.324 GB/s]
saxpy without data transfer: 6.047 ms

CPU

[saxpy ispc]:           [8.482] ms      [35.134] GB/s   [4.716] GFLOPS
[saxpy task ispc]:      [8.202] ms      [36.335] GB/s   [4.877] GFLOPS

可见 saxpy 的 CUDA 版本显著慢于 CPU(ispc) 版本

问题 2

可见计算的部分非常迅速, 瓶颈是数据传输

Part 2

任务 1

首先要确定 线程块的数量和每个线程块中的线程数量

1. 常规设定线程数量为 256
2. 计算所需线程数量
   • 超过 256, <<<threads_num / 256, 256>>>(此处必然可以整除)
   • 不超过 256 <<<1, threads_num>>>

之后就按照 README 文件上的代码翻译即可, 剩余见代码

任务 2

首先 output[idx] = (input[idx] == input[idx + 1]) 再求前缀和

• 此时 output 数组的和就是要返回的值

由于 explicit 的特性

• output[idx] + 1 == output[idx + 1] 的位置就是要找的下标之一
• output[idx] 的值还是答案数组的下标

所以更新答案 result[output[idx]] = idx

Part 3

参考 @kykim0 的 Tile-based 实现

朴素实现

两种并行: 圆 和 像素

渲染顺序的问题可以通过逐像素 遍历每个圆来计算

每帧图像可以拆解为 16 * 16 个 tile, 每个 tile 负责 64 * 64 个像素

每个 tile 可以作为一个线程块, 每个线程块中 1024 个线程负责 1024 个像素

计算的过程可以照抄原本的代码, 只需要看代码处理一下 tile 的细节

优化实现

朴素实现效率很低, 考虑优化遍历圆的过程, 只遍历经过这个 tile 的圆

这个地方可以使用 圆 的并行, 每个圆并行地计算是否经过这个 tile

使用 sharedMemExclusiveScan 可以进行计算

这个时候需要把圆分批计算, 因为如果很多圆同时计算的话会爆显存

设置每批圆 1024 个, 这样每个 tile 中

1. prefixSumInput 存储当前圆是否与当前 tile 相交
2. prefixSumOutput 存储 explicit 的前缀和
3. 在全局开一个 circleFlag 数组
   大小为 sizeof(short) * NUM_CIRCLES_PROC * numTilesX * numTilesY
   因为每个 tile 中最多能相交 1024 个圆 (同一批)
   • 当前 圆 在 circleFlag 中的下标 flagIndex 为: 1024 * tileIdx + sum[circleId]
   • 在 flagIndex 处记录 circleIdx 备用
   • 注意如果当前的 circle 是最后一个则需要记录表示终止的 -1

使用 像素 并行计算每个像素的值: 当 flagValue 为 -1 时就提前结束循环