关于CUDA PTX代码的讨论

_梦魇花葬 2014-05-24 07:10:56

加精

什么是PTX代码呢？类似于高级语言编译器编译出来的汇编指令，PTX代码就是可以在GPU端执行的汇编指令。查看PTX代码可以更加直观的修改kernel代码。
在这里举个查看PTX代码必要性的例子。了解GPU基础架构的小伙伴们都知道，在GPU的存储层次中有个存储位置叫做“local memory”，它是在寄存器数量不足的情况下的一种补充，虽然名字叫做local memory，但是它位于global memory，这就意味这它的速度远远要慢于寄存器，也就是说我们应该避免使用它。那在什么情况下会使用这块显存呢？翻阅了很多资料，只得出了一种结论——“在寄存器数量不足的情况下，编译器会自动将数据分配到local memory”。但是，在如下的代码中，local memory的使用量不为0，也就是说kernel函数中有变量使用了local memory：

__global__ void test(float *left, float *right, float *result, int size)

{

    if(threadIdx.x < size)

    result[threadIdx.x] = left[threadIdx.x] + right[threadIdx.x];

}

究竟是哪个变量使用了PTX代码呢？通过查看，发现变量size被分配到了local memory中，因为是值传递，如果是指针传递，就不会发生将size分配到local memory中的这种情况了。
小伙伴们也来谈谈PTX代码的特殊经历吧~！！

...全文

3516 45 打赏收藏转发到动态举报

写回复

45 条回复

切换为时间正序

请发表友善的回复…

发表回复

Eloudy 2015-03-26

打赏
举报

回复

ｌｉｎｕｘ中，nvcc -keep 或-ptx，然后到项目文件夹中去寻找　　＊．ｐｔｘ文件ｖｓ２０１０中到项目属性中，找　“配置属性”　＝》“ｃｕｄａ　Ｃ／Ｃ＋＋”　中寻找相关选项。比如：ｋｅｅｐ preprocessed files中选择“是” 望能找到。请问学习ptx的文档是哪个比较好呢？

huikong2013 2014-11-22

打赏
举报

回复

请问楼主，怎样才能把编译的C程序转化成PTX呢，用NVCC？

_梦魇花葬 2014-07-15

打赏
举报

回复

好的，没问题，我会分享一些新的东西，谢谢各位小伙伴的支持！~~

yangfeilong1988 2014-07-15

打赏
举报

回复

一起学习，版主可以提供一些新的！~~

_梦魇花葬 2014-07-15

打赏
举报

回复

希望各位小伙伴都能很好的学习到怎么查看PTX代码，有什么不懂的，可以提出来！

_梦魇花葬 2014-07-14

打赏
举报

回复

很正确，PTX代码是非常有用的，大家一起学习！~~

xx0_0dd815 2014-06-23

打赏
举报

回复

什么是PTX代码呢？类似于高级语言编译器编译出来的汇编指令，PTX代码就是可以在GPU端执行的汇编指令。查看PTX代码可以更加直观的修改kernel代码。

qq_16797523 2014-06-23

打赏
举报

回复

什么是PTX代码呢？类似于高级语言编译器编译出来的汇编指令，PTX代码就是可以在GPU端执行的汇编指令。查看PTX代码可以更加直观的修改kernel代码。

2706556716 2014-06-23

打赏
举报

回复

谢谢，学习了

yuea20000 2014-06-23

打赏
举报

回复

谢谢，学习了。。。。

qd73367 2014-06-23

打赏
举报

回复

谢谢分享,xxxxxxxxxxxx

郑卫星 2014-06-23

打赏
举报

回复

innovy 2014-06-20

打赏
举报

回复

好东西，值得学习

ft8266 2014-06-20

打赏
举报

回复

06-11 13:52

qq_16685649 2014-06-19

打赏
举报

回复

谢谢~！受教了

think93723 2014-06-18

打赏
举报

回复

謝謝分享 thanks

oicq007qiaoshi 2014-06-18

打赏
举报

回复

原来学校内学C语言的时候没当回事儿，毕业后后悔了！亡羊补牢，为时不晚。现在继续学它！！

qq_16638613 2014-06-18

打赏
举报

回复

论坛真的不错初学者要努力

qq_16638187 2014-06-18

打赏
举报

回复

好吧，看不懂的

bh779 2014-06-18

打赏
举报

回复

很实用呀

加载更多回复（18）

内联PTX汇编具有如下形式： asm(“instop”:”type_symbolic”(or):”type_symbolic”(ir),..); asm(“instop”::”type_symbolic”(r)); 第二种形式是无输出操作格式，需要使用”::”指示符其中instop是指令操作 type_symbolic是类型指示符（可选如下），分别对应与PTX中的数据类型： “h”

3.3 内存层次(Memory Hierarchy)在GPU上CUDA线程可以访问到的存储资源有很多，每个CUDA线程拥有独立的本地内存(local Memory)；每一个线程块(block)都有其独立的共享内存(shared memory)，共享内存对于线程块中的每个线程都是可见的，它与线程块具有相同的生存时间；同时，还有一片称为全局内存(global memory)的区域对所有的CUDA线程都...

驱动程序 API 在 cuda 动态库（或）中实现，该库在安装设备驱动程序期间复制到系统上。它的所有入口点都以 cu 为前缀。它是一个基于句柄的命令式 API：大多数对象都由不透明的句柄引用，这些句柄可以指定给函数来操作对象。驱动程序 API 中可用的对象汇总在下表中。在调用驱动程序 API 的任何函数之前，必须使用初始化驱动程序 API。然后必须创建一个附加到特定设备的 CUDA 上下文，并使其成为当前调用主机线程，如上下文中所述。在 CUDA 上下文中，内核作为 PTX 或二进制对象由主机代

内联PTX汇编具有如下形式：asm(“instop”:”type_symbolic”(or):”type_symbolic”(ir),..);asm(“instop”::”type_symbolic”(r));第二种形式是无输出操作格式，需要使用”::”指示符其中instop是指令操作type_symbolic是类型指示符（可选如下），分别对应与PTX中的数据类型：“h” .s16, .u16“r” .s32, .u32“l” .s64, .u64“f” .f32“d” .f64例如：//c=a+bfloa

CUDA C 编程

579

社区成员

2,918

社区内容

发帖

与我相关

我的任务

社区管理员

加入社区

近7日
近30日
至今

加载中

查看更多榜单

社区公告

暂无公告

试试用AI创作助手写篇文章吧

+ 用AI写文章