ptx相关内容

使用内联PTX asm()指令时,"volatile"有什么作用?

当我们在通常的C/C ++ CUDA代码中编写内联PTX程序集时,例如: __ device__ __inline__ uint32_t bfind(uint32_t val){uint32_t ret;asm("bfind.u32%0,%1;":"= r"(ret):"r"(val));返回ret} 我们可以在 asm 之后添加 volatile 关键字,例如: __ device_ ..
发布时间:2021-04-27 20:13:33 其他开发

如何找到活跃的SM?

是否可以通过任何方式了解空闲/活动SM的数量?还是至少要读取每个SM的电压/功率或温度值,才能知道其是否正常工作? (在gpu设备上执行某些作业时实时)。 %smid帮助我了解了每个SM的ID。 感谢和问候, Rakesh 解决方案 CUDA分析工具接口( CUPTI )包含事件API可以对GPU PM计数器进行运行时采样。 CUPTI SDK随CUDA Toolkit一起提 ..
发布时间:2020-10-13 00:47:32 其他开发

将内核链接到PTX功能

我可以使用PTX文件中包含的PTX函数作为外部设备函数来将其链接到另一个应调用该函数的.cu文件吗? 这是 CUDA-将内核链接在一起中的另一个问题,其中函数本身不包含在。 cu文件,但我宁愿以某种方式链接PTX功能。 解决方案 您可以自己加载包含PTX代码的文件 cuModuleLoad 和 cuModuleGetFunction 来自文件系统的代码,如下所示: CUmod ..
发布时间:2020-10-13 00:45:49 其他开发

直接将PTX程序传递给CUDA驱动程序

CUDA驱动程序API提供了从文件系统加载包含PTX代码的文件。通常执行以下操作: CUmodule模块; CUfunction函数; const char * module_file =“ my_prg.ptx”; const char * kernel_name =“ vector_add”; err = cuModuleLoad(& module,module ..
发布时间:2020-10-13 00:41:58 其他开发

如何在运行时生成,编译和运行CUDA内核

好吧,我有一个很棘手的问题:) 让我们从我拥有的东西开始: 数据,大​​量数据,已复制到GPU 程序,由CPU(主机)生成,需要对该数组中的每个数据进行评估 程序更改非常频繁,可以生成为CUDA字符串,PTX字符串或其他形式(? ),并且每次更改后都需要重新评估 我想要的是:基本上只是想做出这尽可能有效(快速),例如。避免将CUDA编译为PTX。解决方案甚至可以完全是特定于设 ..
发布时间:2020-10-06 20:40:19 其他开发

CUDA 9中附加了一些以`_sync()`命名的内部函数;语义相同吗?

在CUDA 9中,nVIDIA似乎有了这种“合作团体"的新概念;由于某种原因(我不太清楚),现在不推荐使用__ballot()(= CUDA 9),而推荐使用__ballot_sync().是别名还是语义发生了变化? ...类似的问题,对于现在已在其名称中添加了__sync()的其他内建程序. 解决方案 没有语义是不一样的.函数调用本身是不同的,一个不是另一个的别名,已经公开了新功能 ..
发布时间:2020-07-31 01:53:02 其他开发

对于PTX文件中的某些损坏的名称,c ++ filt不够积极

我正在通过c ++ filt过滤已编译的PTX,但是它只会使某些名称/标签不符合要求,并保留某些原样.例如,这: func (.param .b32 func_retval0) _ZN41_INTERNAL_19_gather_bits_cpp1_ii_56538e7c6__shflEiii( .param .b32 _ZN41_INTERNAL_19_gather_bit ..
发布时间:2020-07-01 04:09:05 其他开发

在一维网格中计算经向ID/车道ID的最有效方法是什么?

在CUDA中,每个线程都知道其在网格中的块索引以及该块内的线程索引.但是似乎没有两个重要的值可供使用: 其索引为经线内的车道(其“车道ID") 它是块内泳道的经线索引(其“经线ID") 假设网格是一维的(又称线性,即blockDim.y和blockDim.z为1),显然可以这样获得: enum : unsigned { warp_size = 32 }; auto lane_i ..
发布时间:2020-05-21 20:50:59 其他开发

Cuda PTX注册声明和使用

我尝试减少内核中使用寄存器的数量,所以我决定尝试内联的PTX。 这个内核: #define Feedback(a,b,c,d,e)d ^ e ^(a& c)^(a& e)^ c和c(a& b)^(a& b& d)^(a& b& d) c) __global__ void测试(unsigned long a,unsigned long b,unsigned long c,unsi ..
发布时间:2017-03-05 19:18:03 其它硬件开发

使用PTX在C ++ / CUDA程序中计算浮点运算的方法

我有一个有点大的CUDA应用程序,我需要计算获得的GFLOPs。 我正在寻找一种简单而且通用的计算浮点运算数的方法。 可以从生成的浮点运算计数PTX代码(如下所示),使用汇编语言中的预定义fpo的列表?基于代码,计数可以通用吗?例如, add.s32%r58,%r8,-2; 计为一个浮点运算? EXAMPLE: BB3_2: .loc 2 108 1 mov.u32%r ..
发布时间:2017-03-05 19:04:10 其它硬件开发

CUDA的内联PTX代码的语法

如Nvidia的Inline PTX Assembly文档中所述,使用内联汇编的语法是: asm(“temp_string”:“constraint”(output):“constraint”(input)) ; 以下是两个示例: asm(“vadd.s32.s32.s32%0,%1.h0 ,%2.h0;“:”= r“(v):”r“(a),”r“(b)); “as”(“vadd ..
发布时间:2017-03-05 15:27:35 其它硬件开发

与CUDA PTX代码和寄存器存储器混淆

:) 虽然我试图管理我的内核资源,我决定看看PTX,但有一些事情,我不明白。这是一个非常简单的内核我写的: __ global__ void foo(float * out,float * in,uint32_t n) { uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x; uint32_t one = 5; out ..
发布时间:2017-03-04 16:03:47 其它硬件开发

我应该看看PTX来优化我的内核吗?如果是,如何?

您是否建议您阅读内核的PTX代码,以便进一步优化内核? 一个例子:我读到,可以从PTX代码找出自动循环展开是否有效。如果不是这种情况,则必须在内核代码中手动展开循环。 PTX代码还有其他用例吗? 你看看你的PTX代码吗? 我在哪里可以找到如何读取CUDA为我的内核生成的PTX代码? 解决方案 关于PTX的第一点是,它只是在GPU上运行的代码的一个中间表示 - 虚拟机汇编语言 ..
发布时间:2017-03-04 15:59:51 其它硬件开发

CUDA / PTX 32位与64位

CUDA编译器具有生成32位或64位PTX的选项。这些之间有什么区别?是否喜欢x86,NVidia GPU实际上有32位和64位ISA? 解决方案 指针肯定是最明显的区别。 64位机器型号启用64位指针。 64位指针支持多种功能,例如大于4GB的地址空间,统一虚拟寻址。统一虚拟寻址又启用其他功能,例如 GPUDirect对等-Peer 。 CUDA IPC API 也取决于64位机器型号。 ..
发布时间:2017-03-04 15:10:46 其它硬件开发

是否可以将汇编指令放入CUDA代码?

我想在CUDA C代码 中使用汇编代码,以减少昂贵的执行 ,因为我们在c编程中使用 asm 。 有可能吗? 解决方案 不,你不能,没有什么像C / C ++的asm结构。您可以做的是调整生成的PTX程序集,然后与CUDA一起使用。 请参阅这个例子。 但是对于GPU来说,程序集优化是不必要的,你应该首先做其他优化,例如存储器合并和占用。请参见 CUDA最佳做法指南有关更 ..
发布时间:2017-03-04 15:09:13 其它硬件开发

CUDA:如何使用-arch和-code和SM vs COMPUTE

我仍然不知道如何正确地指定架构代码生成时使用nvcc构建。我知道有机器代码和PTX代码嵌入在我的二进制,这可以通过控制器开关 -code 和 - arch (或两者的组合使用 -gencode )。 到这个除了两个编译器标志外,还有两个指定架构的方法: sm_XX 和 compute_XX ,其中 compute_XX 指向真实体系结构的虚拟和 sm_XX 。标志 -arch 只使用虚拟架构 ..
发布时间:2017-03-04 13:07:06 其它硬件开发

如何编译PTX代码

我需要修改PTX代码并直接编译。原因是我想要一些特定的指令紧接着彼此,并且很难写一个cuda代码,导致我的目标PTX代码,所以我需要直接修改ptx代码。 问题是我可以编译它(fatbin和cubin),但我不知道如何编译这些(.fatbin和.cubin)到“X.o”文件。 感谢您的帮助。 解决方案 一个方法来执行这个有序的 nvcc 命令,但我不知道它,并没有发现它。 然而 ..
发布时间:2017-03-04 12:58:37 其它硬件开发

在内联ptx中加载函数参数

我有以下函数与内联汇编在32位Visual Studio 2008调试模式下工作正常: __ device__ void add(int * pa,int * pb) { asm(“.reg .u32 s 3;”::); asm(“。reg .u32 r 14&”;“::); asm(“ld.global.b32 s0,[%0];”::“r”(& pa)); //加载pa的 ..
发布时间:2017-03-04 12:24:47 其它硬件开发

检测Thrust变换的ptx内核

我有以下的thrust :: transform调用。 my_functor * f_1 = new my_functor thrust :: transform(data.begin(),data.end(),data.begin(),* f_1); 我想检测它在PTX文件中的相应内核。但是有许多内核在其名称中包含my_functor。 例如 - ..
发布时间:2017-03-04 12:03:35 其它硬件开发