风辰的cuda入门教程

zyren988

赞一个

ali_m

这个和论坛里置顶的那个有区别吗?

悠闲的小猫

热烈庆祝风辰老大的又一大做问世！！悠猫前来捧场！！

yyfn风辰

悠闲的小猫

老大，您的文章里面如果有错误，可以挑刺吗？额。。。包括错别字？

悠闲的小猫

原帖由 yyfn风辰 于 2010-11-16 10:51 发表
多给点意见啊

风大，
您看这里：
（1）
第三页倒数第二段，CUDA compiler是host c/c++ compiler的前端才对。而且。对于在windows下，nvcc并不需要VS, 而只要基本的那个microsoft optimization c++ compiler就可以工作了(cl)。可以单独从windows sdk中安装。而且是免费的。
(2)还是这一段，栅栏同步您指的是？

(3)第12页，您说在tesla(1.x)架构上，如果有shared memory bank conflicts, 那么最严重的时候甚至会比访问global memory还慢？假设该16个线程都读取了同一bank的不同位置, 那么他们需要16次读取。在1.x上，在只读取shared memory的情况下，1次需要1个clock，即1.x读取shared和读取寄存器的速度是一样的, 那么最严重的时候需要16个clocks。然后考虑到即使是完全合并的对global memory的读取(最快)，一次也需要400-600个clocks。这样，即使是在读取的时候，因为最严重的bank conflicts的情况下，读取shared memory所需要的周期也只global memory的3%~4%.我觉得您说的是不恰当的。

(4)倒数第二段，fermi的对shared memory的读取是以warp位单位的(一次32个)，而不再是以half-warp(16个)位单位了。

(5)第18页的图下面有段文字。

(6)第18页倒数第二段。您的延迟指的是什么？数据从哪里来的？
原文“每次运算需要4个时钟周期(SP周期)”，这里有很大的误导性！很容易让人以为，一个运算N次，一共要4*N的cuda core clocks.
这段还有一些别的小问题，以后讨论。

以上是对您的文章的第一章说的。

蓝云风翼

学习ing

yyfn风辰

一、说host C++是前端是因为，语法什么基本上都是ｈｏｓｔ处理的
二、指的是barrier, __syncthreads()
三、shared有四十左右个clock
四、手册上是这样的，但是实际上每次最多还是１６个
六，因为以前是８个ｓｐ，而实际上一个ｗａｒｐ有３２线程，ｓｐ要连续执行４个周期才能给你个ｔｈｒｅａｄ一个周期

悠闲的小猫

原帖由 yyfn风辰 于 2010-12-2 12:43 发表
一、说host C++是前端是因为，语法什么基本上都是ｈｏｓｔ处理的

嗯.一般我们说前端, 后端, 是按照编译流程顺序来的. 比如说,  gcc可以使用llvm作为后端, 指的就是gcc编译->LLVM机器码->X86机器码.
那么由于, 在交给nvcc配合gcc. 是nvcc先整理.cu, 提出device code部分, 剩下的交给gcc, 那么是不是可以认为是nvcc是gcc的前端呢?

三、shared有四十左右个clock

据我所掌握的资料, G80/200系列的机器码(注意是机器码, 不是ptx), 对于shared memory是直接访问的(虽然在ptx可以看到ld.shard之类的语句, 但是会被jit成一条直接使用指定shared memory单元的指令.). 而且据说在1.x上, shared就是用(看不见的)registers实现的. 而且很多资料显示, 在1.x上, 读取shared只需要1个周期(比2.x的shared都快!). 而不是需要40多个.

四、手册上是这样的，但是实际上每次最多还是１６个.

何以见得呢? 如果真的是16个. 那岂不是手册上一再强调的在2.x上, 要注意32个bank情况下, 导致的1.x不会有conflicts, 而2.x会有的情况是不必要的了? 如果不必要. 为了要额外强调? 您有详细资料证明是16个吗?

每次运算需要4个时钟周期(SP周期), 包括取出数据, 计算, 写回.

六，因为以前是８个ｓｐ，而实际上一个ｗａｒｐ有３２线程，ｓｐ要连续执行４个周期才能给你个ｔｈｒｅａｄ一个周期

是不是这两句话有冲突?

yyfn风辰

一、可能真的是我说错了，呵呵！以后不用这种说法了
三和四是我自己测试过的，在GTX 295和C2050上读取一个shared数据38-42个clock, 而在这两种机器上读取一个warp的shared时间差不多

六，我以后会注意说法，这个主要是受别人的书的影响。