 C++高性能并行编程与优化 -  课件 - 08 CUDA 开启的 GPU 编程如果你试图分离声明和定义,调用另一个文件里 的 __device__ 或 __global__ 函数,就会出错 。 分离 __device__ 函数的声明和定义:解决 • 开启 CMAKE_CUDA_SEPARABLE_COMPILATION 选 项(设为 ON ),即可启用分离声明和定义的支持。 • 不过我还是建议把要相互调用的 __device__ 函数放在 同一个文件,这样方便编译器自动内联优化(第四课讲 glDispatchComputeIndirect 的 API 和这个很像,但毕竟没有 CUDA 可以直接在核函数里调用核函数并指定参数这么方便…… 不过,这个功能同样需要开启 CUDA_SEPARABLE_COMPILATION 。 第 2 章:内存管理 如何从核函数里返回数据? • 我们试着把 kernel 的返回类型声明为 int ,试 图从 GPU 返回数据到 CPU 。 • 但发现这样做会在编译期出错,为什么?0 码力 | 142 页 | 13.52 MB | 1 年前3 C++高性能并行编程与优化 -  课件 - 08 CUDA 开启的 GPU 编程如果你试图分离声明和定义,调用另一个文件里 的 __device__ 或 __global__ 函数,就会出错 。 分离 __device__ 函数的声明和定义:解决 • 开启 CMAKE_CUDA_SEPARABLE_COMPILATION 选 项(设为 ON ),即可启用分离声明和定义的支持。 • 不过我还是建议把要相互调用的 __device__ 函数放在 同一个文件,这样方便编译器自动内联优化(第四课讲 glDispatchComputeIndirect 的 API 和这个很像,但毕竟没有 CUDA 可以直接在核函数里调用核函数并指定参数这么方便…… 不过,这个功能同样需要开启 CUDA_SEPARABLE_COMPILATION 。 第 2 章:内存管理 如何从核函数里返回数据? • 我们试着把 kernel 的返回类型声明为 int ,试 图从 GPU 返回数据到 CPU 。 • 但发现这样做会在编译期出错,为什么?0 码力 | 142 页 | 13.52 MB | 1 年前3
 C++高性能并行编程与优化 -  课件 - 07 深入浅出访存优化波操作,就属于插桩。有的 插桩内核各轴向是对称的(比如高斯模糊),有的是单单往一个方向延伸 很长(比如径向模糊),有的内核是正方形(箱滤波)。 • 人工智障圈子里好像管这个叫卷积( convolution )什么的。 X 方向插桩:手动预取下一缓存行 • X 方向的插桩,因为最内层有另一个长度仅为 nblur*2+1 的顺序读取,让 CPU 误以为 t++ 是我 们想要的顺序读取,然而0 码力 | 147 页 | 18.88 MB | 1 年前3 C++高性能并行编程与优化 -  课件 - 07 深入浅出访存优化波操作,就属于插桩。有的 插桩内核各轴向是对称的(比如高斯模糊),有的是单单往一个方向延伸 很长(比如径向模糊),有的内核是正方形(箱滤波)。 • 人工智障圈子里好像管这个叫卷积( convolution )什么的。 X 方向插桩:手动预取下一缓存行 • X 方向的插桩,因为最内层有另一个长度仅为 nblur*2+1 的顺序读取,让 CPU 误以为 t++ 是我 们想要的顺序读取,然而0 码力 | 147 页 | 18.88 MB | 1 年前3
共 2 条
- 1













