基于CUDA的GPU并行计算.pptVIP

  1. 1、有哪些信誉好的足球投注网站(book118)网站文档一经付费(服务费),不意味着购买了该文档的版权,仅供个人/单位学习、研究之用,不得用于商业用途,未经授权,严禁复制、发行、汇编、翻译或者网络传播等,侵权必究。。
  2. 2、本站所有内容均由合作方或网友上传,本站不对文档的完整性、权威性及其观点立场正确性做任何保证或承诺!文档内容仅供研究参考,付费前请自行鉴别。如您付费,意味着您自己接受本站规则且自行承担风险,本站不退款、不进行额外附加服务;查看《如何避免下载的几个坑》。如果您已付费下载过本站文档,您可以点击 这里二次下载
  3. 3、如文档侵犯商业秘密、侵犯著作权、侵犯人身权等,请点击“版权申诉”(推荐),也可以打举报电话:400-050-0827(电话支持时间:9:00-18:30)。
  4. 4、该文档为VIP文档,如果想要下载,成为VIP会员后,下载免费。
  5. 5、成为VIP后,下载本文档将扣除1次下载权益。下载后,不支持退款、换文档。如有疑问请联系我们
  6. 6、成为VIP后,您将拥有八大权益,权益包括:VIP文档下载权益、阅读免打扰、文档格式转换、高级专利检索、专属身份标志、高级客服、多端互通、版权登记。
  7. 7、VIP文档为合作方或网友上传,每下载1次, 网站将根据用户上传文档的质量评分、类型等,对文档贡献者给予高额补贴、流量扶持。如果你也想贡献VIP文档。上传文档
查看更多
基于CUDA的GPU并行计算.ppt

优化原则:active block 一个SM中可以有多个block等待处理,在一个warp需要访问存储器或者同步时,另外一个warp可以使用执行单元的资源 增加active block对提高SM利用率有好处 增加active block只是手段,不是最终的评价标准。最终目的是要隐藏延迟 优化原则:active block 每个SM最多可以有768(G8x,G9x)或者1024(GT200)个active thread 这些active thread最多可以属于8个block 还有受到SM中shared memory和register的制约 最后的active block数量是由以上四个条件中的“短板”决定 并行计算 Vector Reduction with Branch Divergence 0 1 2 3 4 5 7 6 10 9 8 11 0+1 2+3 4+5 6+7 10+11 8+9 0...3 4..7 8..11 0..7 8..15 1 2 3 Thread 0 Thread 8 Thread 2 Thread 4 Thread 6 Thread 10 A simple implementation Assume we have already loaded array into __shared__ float partialSum[] unsigned int t = threadIdx.x; for (unsigned int stride = 1; stride blockDim.x; stride *= 2) { __syncthreads(); if (t % (2*stride) == 0) partialSum[t] += partialSum[t+stride]; } No Divergence until 16 sub-sums Thread 0 0 1 2 3 … 13 15 14 18 17 16 19 0+16 15+31 1 3 4 A better implementation Assume we have already loaded array into __shared__ float partialSum[] unsigned int t = threadIdx.x; for (unsigned int stride = blockDim.x; stride 1; stride 1) { __syncthreads(); if (t stride) partialSum[t] += partialSum[t+stride]; } Memory Coalescing When accessing global memory, peak performance utilization occurs when all threads in a Warp access continuous memory locations. Md Nd W I D T H WIDTH Thread 1 Thread 2 No coalesced coalesced Memory Access Pattern Md Nd W I D T H WIDTH Md Nd Original Access Pattern Tiled Access Pattern Copy into scratchpad memory Perform multiplication with scratchpad values * * * 南京大学多媒体研究所 南京大学多媒体研究所 基于CUDA的GPU并行计算 Compute Unified Device Architecture * 南京大学多媒体研究所 * 内容 概述 CUDA的硬件 CUDA的软件 性能提升 并行计算 概述 摩尔定律 * 南京大学多媒体研究所 * 摩尔定律的失效 晶体管尺寸 流水线级数 能量墙 存储墙 * 南京大学多媒体研究所 * 新摩尔定律 未来计算机硬件不会更快,但会更“宽” ? No longer get faster, just wider 必须重新设计算法 * 南京大学多媒体研究所 * 1 Based on slide 7 of S. Green, “GPU Physics,” SIGGRAPH 2007 GPGPU Course. /s2007/slides/15-GPGPU-physics.pdf What is driving the many-cores? GPU与CPU的不同 The GPU is specialized for compute-

文档评论(0)

docinppt + 关注
实名认证
文档贡献者

该用户很懒,什么也没介绍

1亿VIP精品文档

相关文档