我最近了解了NVCC如何为不同的计算架构编译CUDA设备代码.
根据我的理解,当使用NVCC的-gencode选项时,"arch"是程序员应用程序所需的最小计算体系结构,也是NVCC的JIT编译器将编译PTX代码的最小设备计算体系结构.
我也明白-gencode的"code"参数是NVCC完全编译应用程序的计算架构,因此不需要JIT编译.
在检查了各种CUDA项目Makefile之后,我注意到以下情况经常发生:
-gencode arch=compute_20,code=sm_20
-gencode arch=compute_20,code=sm_21
-gencode arch=compute_21,code=sm_21
经过一些阅读,我发现可以在一个二进制文件中编译多个设备架构 - 在本例中为sm_20,sm_21.
我的问题是为什么需要这么多的arch/code对？以上是否使用了"拱"的所有值？
它之间的区别是什么？
-arch compute_20
-code sm_20
-code sm_21
是自动选择"拱形"字段中最早的虚拟架构,还是存在其他一些模糊行为？
我应该注意其他任何编译和运行时行为吗？
我已经阅读了手册http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#gpu-compilation,我仍然不清楚编译或运行时会发生什么.
在使用nvcc构建时,我仍然不确定如何正确指定代码生成的体系结构.我知道我的二进制文件中嵌入了机器代码和PTX代码,这可以通过控制器开关-code和-arch(或两者的结合使用-gencode)来控制.
现在,根据本除了两种编译器标志也有指定架构方法有两种:sm_XX和compute_XX,其中compute_XX指的是虚拟和sm_XX一个真正的架构.该标志-arch仅采用虚拟体系结构的标识符(例如compute_XX),而-code标志采用真实和虚拟体系结构的标识符.
该文档指出了-arch指定为其编译输入文件的虚拟体系结构.但是,此PTX代码不会自动编译为机器代码,但这是一个"预处理步骤".
现在,-code应该指定汇编和优化PTX代码的架构.
但是,不清楚哪个PTX或二进制代码将嵌入二进制文件中.例如-arch=compute_30 -code=sm_52,如果我指定,这是否意味着我的代码将首先被编译为功能级别3.0 PTX,之后将创建功能级别5.2的机器代码？什么将被嵌入？
如果我只是说明-code=sm_52会发生什么呢？只嵌入V5.2的机器代码是用V5.2 PTX代码创建的？那有什么区别-code=compute_52？
在CUDA 2.0设备上有没有办法只为一个特定变量禁用L1缓存？我知道,一个可以在编译时禁用L1缓存添加标志-Xptxas -dlcm=cg,以nvcc对所有内存操作.但是,我想仅对特定全局变量的内存读取禁用缓存,以便所有其余内存读取都通过L1缓存.
基于我在网络上进行的搜索,可能的解决方案是通过PTX汇编代码.
通过CUDA 5.0编程指南阅读时,我偶然发现了一个名为"漏斗移位"的功能,该功能存在于3.5计算设备中,但不是3.0.它包含注释"参见参考手册",但是当我在手册中搜索"漏斗转换"术语时,我找不到任何内容.
我试着谷歌搜索它,但只在http://www.cudahandbook.com上找到了提及,在第8章:
8.2.3漏斗转换(SM 3.5)
GK110添加了一个64位"漏斗移位"指令,可以通过以下内在函数访问:
__funnelshift_lc():返回左漏斗移位的最重要的32位.
__funnelshift_rc():返回右漏斗移位的最低有效32位.
这些内在函数在sm_35_intrinsics.h中实现为内联设备函数(使用内联PTX汇编程序).
......但它仍然没有解释"左漏斗转移"或"正确的漏斗转移"是什么.
那么,它是什么,需要它在哪里？
有没有人知道如何使用新的LLVM后端使用C/C++代码注释PTX汇编程序？
可以使用CUDA 4.0或更早版本轻松获取它,但在将CUDA工具包升级到版本4.2后,NVCC拒绝所有标记.
我正在研究PTX,我不明白CTA(计算线程数组)与CUDA块的不同之处.
它们是一样的吗？在我看来,现在(我只是在PTX文件的开头)他们是一样的
据我所知,Fermi GPU支持预取L1或L2缓存.但是,在CUDA参考手册中我找不到任何关于它的东西.
Dues CUDA允许我的内核代码预取特定数据到特定级别的缓存？
编辑:这个问题是原版的重新版本,因此前几个回复可能不再相关.
我很好奇设备函数调用强制非内联对设备函数内同步的影响.我有一个简单的测试内核,用于说明相关行为.
内核获取一个缓冲区并将其传递给设备函数,以及共享缓冲区和指示符变量,该变量将单个线程标识为"boss"线程.设备功能有不同的代码:boss线程首先花时间对共享缓冲区执行简单的操作,然后写入全局缓冲区.在同步调用之后,所有线程都写入全局缓冲区.在内核调用之后,主机打印全局缓冲区的内容.这是代码:
CUDA代码:
test_main.cu
#include<cutil_inline.h>
#include "test_kernel.cu"
int main()
{
int scratchBufferLength = 100;
int *scratchBuffer;
int *d_scratchBuffer;
int b = 1;
int t = 64;
// copy scratch buffer to device
scratchBuffer = (int *)calloc(scratchBufferLength,sizeof(int));
cutilSafeCall( cudaMalloc(&d_scratchBuffer,
sizeof(int) * scratchBufferLength) );
cutilSafeCall( cudaMemcpy(d_scratchBuffer, scratchBuffer,
sizeof(int)*scratchBufferLength, cudaMemcpyHostToDevice) );
// kernel call
testKernel<<<b, t>>>(d_scratchBuffer);
cudaThreadSynchronize();
// copy data back to host
cutilSafeCall( cudaMemcpy(scratchBuffer, d_scratchBuffer,
sizeof(int) * scratchBufferLength, cudaMemcpyDeviceToHost) );
// print results
printf("Scratch buffer contents: \t");
for(int i=0; i < …我需要修改PTX代码并直接编译.原因是我想要一个接一个地有一些特定的指令,并且很难编写一个导致我的目标PTX代码的cuda代码,所以我需要直接修改ptx代码.问题是我可以将它编译为(fatbin和cubin),但我不知道如何将这些(.fatbin和.cubin)编译为"Xo"文件.
我想在CUDA C代码中使用汇编代码,以减少昂贵的执行,因为我们在c编程中使用asm.
可能吗？