标签: ptx

我想在CUDA C代码中使用汇编代码,以减少昂贵的执行,因为我们在c编程中使用asm.

可能吗？

您是否建议阅读内核的PTX代码以进一步优化内核？

一个例子:我读过,如果自动循环展开有效,可以从PTX代码中找到.如果不是这种情况,则必须在内核代码中手动展开循环.

• PTX代码还有其他用例吗？
• 你看看你的PTX代码了吗？
• 在哪里可以找到如何读取CUDA为我的内核生成的PTX代码？

...仅在PTX手册中提到过.关于它们有什么用处或如何使用它们没有任何暗示.

有谁知道更多？我只是错过了一个共同的概念吗？

我正在尝试创建结构类型的LLVM值.我正在使用LLVM-C接口并找到一个功能:

LLVMValueRef    LLVMConstStruct (LLVMValueRef *ConstantVals, unsigned Count, LLVMBool Packed)

如果所有成员都是由LLVMConstXXX()创建的常量值,这将正常工作,它将生成如下代码:

store { i32, i32, i32 } { i32 1, i32 2, i32 3 }, { i32, i32, i32 }* %17, align 4

但问题是如果成员不是常数,它将产生如下内容:

%0 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
store { i32, i32, i32 } { i32 1, i32 %0, i32 3 }, { i32, i32, i32 }* %17, align 4

当我将这段LLVM代码发送到NVVM(Nvidia PTX后端)时,它说:

模块0(27,39):解析错误:无效使用函数本地名称

所以,我不知道这个结构值创建是否正确.我需要的是一个值,而不是分配的内存.

有人有想法吗？

问候,翔.

CUDA驱动程序API提供从文件系统加载包含PTX代码的文件.通常会做以下事情:

CUmodule module;
CUfunction function;

const char* module_file = "my_prg.ptx";
const char* kernel_name = "vector_add";

err = cuModuleLoad(&module, module_file);
err = cuModuleGetFunction(&function, module, kernel_name);

如果在运行时(运行中)生成PTX文件,则通过文件IO似乎是浪费(因为驱动程序必须再次加载它).

有没有办法直接将PTX程序传递给CUDA驱动程序(例如作为C字符串)？

考虑这3个简单的最小内核.他们的注册用量远高于我的预期.为什么？

A:

__global__ void Kernel_A()
{  
//empty
}

对应的ptx:

ptxas info    : Compiling entry function '_Z8Kernel_Av' for 'sm_20'
ptxas info    : Function properties for _Z8Kernel_Av
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 2 registers, 32 bytes cmem[0]

B:

template<uchar effective_bank_width>
__global__ void  Kernel_B()
{
//empty
}

template
__global__ void  Kernel_B<1>();

对应的ptx:

ptxas info    : Compiling entry function '_Z8Kernel_BILh1EEvv' for 'sm_20'
ptxas info    : Function properties for _Z8Kernel_BILh1EEvv
    0 bytes …

好吧,我有一个非常微妙的问题:)

让我们从我拥有的东西开始:

1. 数据,大量数据,复制到GPU
2. 由CPU(主机)生成的程序,需要针对该阵列中的每个数据进行评估
3. 该程序的变化很频繁,可以作为CUDA字符串,字符串PTX或别的东西(？)产生并需要重新评估每一个变化之后

我想要的:基本上只是想让它尽可能有效(快速),例如.避免将CUDA编译为PTX.解决方案甚至可以完全针对特定设备,这里不需要大的兼容性:)

我所知道的:我已经知道函数cuLoadModule,它可以从存储在文件中的PTX代码加载和创建内核.但我想,必须有一些其他方法可以直接创建内核,而不必先将其保存到文件中.或者也许可以将其存储为字节码？

我的问题:你会怎么做？您可以发布一个示例或链接到类似主题的网站吗？TY

编辑:好了,PTX内核可以直接从PTX字符串(char数组)运行.无论如何我仍然想知道,有没有更好/更快的解决方案呢？仍然存在从字符串到某些PTX字节码的转换,应该可以避免.我也怀疑,从PTX创建设备特定的Cuda二进制文件的一些聪明的方法可能存在,这将删除JIT编译器滞后(很小,但如果你有大量的内核要运行它可以加起来):)

CUDA编译器可以选择生成32位或64位PTX.这些有什么区别？对于x86来说,NVidia GPU实际上有32位和64位ISA吗？或者它只与主机代码有关？

在CUDA中,每个线程都知道它在网格中的块索引和块内的线程索引.但似乎没有明确可用的两个重要值:

• 它的指数作为其经线内的一条车道(其"车道ID")
• 经线的索引,它是块内的一个通道(它的"warp id")

假设网格是一维(又名线性的,即blockDim.y和blockDim.z是1),可以明显地获得这些如下:

enum : unsigned { warp_size = 32 };
auto lane_id = threadIdx.x % warp_size;
auto warp_id = threadIdx.x / warp_size;

如果您不信任编译器来优化它,您可以将其重写为:

enum : unsigned { warp_size = 32, log_warp_size = 5 };
auto lane_id = threadIdx.x & (warp_size - 1);
auto warp_id = threadIdx.x >> log_warp_size;

这是最有效的事情吗？对于每个线程来说,计算它仍然需要很多浪费.

_{(受这个问题的启发.)}

对于我的项目,我正在以两种不同的方式为某些功能生成PTX指令.第一种方法使用CUDA C来实现函数,使用nvcc来编译它们 nvcc -ptx <file>.cu -o <file>.ptx.另一种方法用不同的语言编写代码,从中生成LLVM IR,并使用NVPTX后端将其编译为ptx.我在这里遇到的问题是某些功能在第二种情况下表现更差.其他功能导致或多或少的可比性能.

现在我想知道为什么某些函数的性能存在这样的差异(以及为什么其他函数没有),但是使用nsight进行分析并没有给我任何好的想法.

我发现的唯一区别是寄存器使用情况.在生成的ptx代码中,我可以看到以下内容:

使用nvcc编译

.reg .u32 %r<8>;
.reg .u64 %rd<17>;
.reg .f32 %f<8>;
.reg .pred %p<5>;

使用nvptx编译

.reg .pred %p<396>;
.reg .s16 %rc<396>;
.reg .s16 %rs<396>;
.reg .s32 %r<396>;
.reg .s64 %rl<396>;
.reg .f64 %fl<396>;

据我所知,这表示使用的虚拟寄存器的数量和类型,但正如您可以清楚地看到的,在第二种情况下这是不正确的.在使用nsight进行性能分析后,我可以看到实际使用的寄存器/线程的数量在第一种情况下为8,在第二种情况下为31.当然,这可能表明为什么第二种情况下的代码速度较慢,但​​问题是我使用NVPTX从LLVM IR编译到ptx的所有函数都存在这个问题.它们都有396个使用过的虚拟寄存器和nsight报告31个所有寄存器/线程,尽管某些函数产生的性能几乎与第一种情况完全相同.

这个注册是我减速的问题吗？为什么它不影响所有功能？如果不是,可能导致经济放缓的原因是什么？你能给出我应该看的方向的任何提示吗？

谢谢!

(使用的LLVM版本是3.3)

编辑:我注意到的另一个不同之处是失速原因:

NVCC:

NVPTX:

显然,"其他"原因相对增加.也许这可以解释这个问题？

编辑:添加了ptx源代码

此处显示的函数将数据从全局内存复制到共享内存.然后每个线程将其自己的元素和前一个元素与数组中的最后一个元素进行比较.如果比较为正,则将索引写入输出数组.

1)使用NVPTX将LLVM IR编译为PTX

// .globl   julia_cuda_find_weighted_median18585
.entry julia_cuda_find_weighted_median18585(
    .param .u64 .ptr .global .align 4 julia_cuda_find_weighted_median18585_param_0,
    .param .u64 .ptr .global .align 4 julia_cuda_find_weighted_median18585_param_1
)                                       // @julia_cuda_find_weighted_median18585
{
    .reg …