当使用内联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)); return ret; } 我们可以在asm之后添加volatile关键字,例如: __device__ __inline__ uint32_t bfind(uint32_t val) { uint32_t ret; asm volatile ("bfind.u32 %0, %1;" : "=r"(ret): "r"(val)); return ret; } cuda 编译器假定asm()语句没有副作用,除了更改输出操作数.为了确保在生成PTX期间未删除或移动asm,您应该使用挥发性关键字 我不明白这意味着什么.所
2 2023-11-25
编程技术问答社区
cuda:设备函数内联和不同的.cu文件
两个事实:CUDA 5.0可让您在不同对象文件中编译CUDA代码,以便以后链接. CUDA体系结构2.x不再自动发挥作用. 像往常一样,在C/C ++中,我在functions.cu中实现了一个函数__device__ int foo(),并将其标头放在functions.hu中.函数foo在其他CUDA源文件中调用. 当我检查functions.ptx时,我看到foo()溢出到本地内存.出于测试目的,我评论了foo()的所有肉,然后根据.ptx将其散布在本地记忆中. (我无法想象它是什么,因为该功能无助!) 但是,当我将foo()的实现移动到标题文件functions.hu并添加__forceinline__限定符时,则没有写入本地内存! 这里发生了什么?为什么CUDA不在围困这样简单的功能? 单独的标头和实现文件的全部要点是使我的生活更轻松地维护代码.但是,如果我必须在标题和__forceinline__中粘贴一堆功能(或所有功能),那么它会失败CUDA
74 2023-05-23
编程技术问答社区
在一维网格中计算经线ID/车道ID的最有效方法是什么?
在CUDA中,每个线程都知道其网格中的块索引,并知道块内的螺纹索引.但是两个重要的值似乎没有明确可用: 其索引是其经线内的车道(其"车道ID") 它在块内(其"扭曲ID")内的纱的索引 假设网格是1维的(又称线性,即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
16 2023-05-12
编程技术问答社区
CUDA PTX代码 %envreg<32>特殊寄存器
我尝试运行一个由.cl内核生成的PTX汇编代码,并使用CUDA驱动程序API生成.我采取的步骤是这些(标准OPENCL程序): 1)加载.cl内核 2)JIT编译 3)获取编译的PTX代码并保存. 到目前为止一切都很好. 我注意到PTX汇编内部的一些特殊寄存器,%envreg3,%envreg6等用驱动程序API代码.因此,该代码落入无限循环中,并且无法进行旋转.但是,如果我手动设置值(nore恰好是我将%envreg6替换为ptx中的块),则代码正在执行,并且我获得了正确的结果(与CPU结果相比正确). > 有人知道我们如何设置值在这些寄存器上,还是我缺少某些东西的话?即Culaunchkernel上的标志,将值设置为这些寄存器? 解决方案 您正在尝试编译OpenCL内核,并使用CUDA驱动程序API运行它. NVIDIA驱动程序/编译器接口在OpenCL和CUDA之间有所不同,因此您不支持您要做的事情,并且从根本上无法使用. 大概,
54 2023-05-09
编程技术问答社区
如何找到活跃的SM?
有什么方法可以知道免费/活动SMS的数量?或至少要读取每个SM的电压/功率或温度值,我可以知道它是否有效? (实时在GPU设备上执行某些作业时). %smid帮助我了解了每个SM的ID.类似的东西会有所帮助. 谢谢,问候, rakesh 解决方案 CUDA分析工具接口( smpling事件 . 以下一个或多个计数器将为您提供SM活动的好主意: active_cycles :多处理器的循环数量至少具有一个主动纱. Active_warps :每个周期的活动扭曲的累积数量.对于每个周期,它会通过周期中的主动扭曲数量增加 可以在0到{48,64}的范围内.
16 2023-05-02
编程技术问答社区
线程数对bar.arrive PTX障碍物同步指令的意义是什么?
它被提及bar.sync a{, b}; bar.arrive a, b; 其中 源操作数A指定逻辑屏障资源为立即常数或具有值0到15的登记.操作数B指定参与屏障的线程数. 它还显示了一个示例,其中使用这些说明建立了生产者消费者模型: // Producer code places produced value in shared memory. st.shared [r0],r1; bar.arrive 0,64; ... // Consumer code, reads value from shared memory bar.sync 0,64; ld.shared r1,[r0]; ... 我没有完全获得bar.arrive中操作数b的目的.虽然bar.sync中的此类操作数可用于控制屏障中涉及的线程数并等到达到螺纹计数,但我对bar.arrive的使用对我来说尚不清楚. 解决方案 当所有线程到达障碍时,都会发生两件事: 所有等待线
28 2023-04-13
编程技术问答社区
我如何让PTX文件执行
我知道如何从.cu生成.ptx文件,以及如何从.ptx.生成.cubin文件,但我不知道如何获得最终可执行文件. 更具体地说,我有一个sample.cu文件,该文件已编译为sample.ptx.然后,我使用NVCC将sample.ptx编译为sample.cubin.但是,如果没有主机代码,就无法直接执行此.cubin文件.如何将文件链接到我的原始.cu文件以生成最终可执行文件? 解决方案 您应该能够使用CumoduleLeLoLoDdataex直接从CUDA驱动程序API运行PTX代码.有一个示例其他解决方案 在CUDA 4.0开始,您可以在CUDA内核中写入inline ptx.
18 2023-04-13
编程技术问答社区
如何在运行时生成、编译和运行CUDA内核
好吧,我有一个非常微妙的问题:) 让我们从我拥有的东西开始: 数据,大​​量数据,复制到gpu 程序,由CPU(主机)生成,需要评估该数组中的每个数据 程序更改非常频繁地,可以作为cuda字符串,ptx字符串或其他内容(?)生成(?),并且在每次更改之后需要重新评估 li> 我想要的:基本上只想使其尽可能有效(快速),例如.避免将CUDA汇编为PTX.解决方案甚至可以完全是特定于设备的,这里不需要大的兼容性:) 我所知道的:我已经知道功能 culoadModule ,可以从文件中存储的PTX代码加载和创建内核.但是我认为,必须有其他方法直接创建内核,而无需将其保存为首先文件.也许可以将其作为字节码存储? 我的问题:您将如何做?您可以发布一个带有类似主题的网站的示例或链接吗? ty 编辑:好的,现在,ptx内核可以是从PTX String(char数组).无论如何,我仍然想知道,是否有一些更好/更快的解决方案?仍然存在从字符串到某些PTX字节码的转换
92 2023-03-22
编程技术问答社区
CUDA仅对一个变量禁用L1高速缓存
在CUDA 2.0设备上有什么办法可以禁用一个特定变量的L1高速缓存? 我知道,可以在编译时间禁用L1缓存,将标志-Xptxas -dlcm=cg添加到nvcc中,以供所有内存操作. 但是,我只想在特定的全局变量上禁用缓存以读取内存,以便所有其余内存都读取要通过L1缓存. 基于我在网络中进行的搜索,可能的解决方案是通过PTX汇编代码. 解决方案 如上所述,您可以使用内联PTX,这是一个示例: __device__ __inline__ double ld_gbl_cg(const double *addr) { double return_value; asm("ld.global.cg.f64 %0, [%1];" : "=d"(return_value) : "l"(addr)); return return_value; } 您可以通过将.f64换成.f32(float)或.s32(int)等,可以轻松地改变此问题) etc et et et e
128 2023-03-03
编程技术问答社区
Cuda PTX寄存器的声明和使用
我正在尝试减少内核中使用寄存器的数量,因此我决定尝试Inline PTX. 此内核: #define Feedback(a, b, c, d, e) d^e^(a&c)^(a&e)^(b&c)^(b&e)^(c&d)^(d&e)^(a&d&e)^(a&c&e)^(a&b&d)^(a&b&c) __global__ void Test(unsigned long a, unsigned long b, unsigned long c, unsigned long d, unsigned long e, unsigned long f, unsigned long j, unsigned long h, unsigned long* res) { res[0] = Feedback( a, b, c, d, e ); res[1] = Feedback( b, c, d, e, f ); res[2] = Feedback( c, d, e, f,
24 2023-03-02
编程技术问答社区
一种使用PTX计算C++/CUDA程序中浮点运算的方法
我有一个较大的CUDA应用程序,我需要计算已达到的Gflops. 我正在寻找一种简单的,也许是计算浮点操作数量的一种通用方法. 使用汇编语言中的预定义FPO列表,可以从生成的PTX代码(如下所示)计数浮点操作吗?基于代码,可以使计数成为通用吗?例如,add.s32 %r58, %r8, -2;是否将其视为一个浮点操作? 示例: BB3_2: .loc 2 108 1 mov.u32 %r8, %r79; setp.ge.s32 %p1, %r78, %r16; setp.lt.s32 %p2, %r78, 0; or.pred %p3, %p2, %p1; @%p3 bra BB3_5; add.s32 %r58, %r8, -2; setp.lt.s32 %p4, %r58, 0; setp.ge.s32 %p5, %r58, %r15; or.pred %
24 2023-03-02
编程技术问答社区
c++filt对于PTX文件中的一些杂乱无章的名称不够积极。
我正在通过C ++ Filt过滤我的PTX,但它只会删除某些名称/标签,并留下一些AS-IS.例如,此: func (.param .b32 func_retval0) _ZN41_INTERNAL_19_gather_bits_cpp1_ii_56538e7c6__shflEiii( .param .b32 _ZN41_INTERNAL_19_gather_bits_cpp1_ii_56538e7c6__shflEiii_param_0, .param .b32 _ZN41_INTERNAL_19_gather_bits_cpp1_ii_56538e7c6__shflEiii_param_1, .param .b32 _ZN41_INTERNAL_19_gather_bits_cpp1_ii_56538e7c6__shflEiii_param_2 ) 被删除了: .func (.param .b32 func_retv
24 2023-03-01
编程技术问答社区
将一个内核与一个PTX函数连接起来
我可以使用PTX文件中包含的PTX函数作为外部设备函数将其链接到另一个.cu文件,该文件应调用该函数? 这是 cuda-链接内核功能本身在哪里, .cu文件中未包含,但我宁愿以某种方式链接PTX函数. 解决方案 您可以从cuModuleLoad和cuModuleGetFunction中加载在您自己的代码中包含PTX代码的文件,如下所示: CUmodule module; CUfunction function; const char* module_file = "my_ptx_file.ptx"; const char* kernel_name = "my_kernel_name"; err = cuModuleLoad(&module, module_file); err = cuModuleGetFunction(&function, module, kernel_name); 您也可以将PTX代码直接作为字符串传递给CUDA驱动程序,请参见.
28 2023-02-25
编程技术问答社区
有可能把汇编指令放到CUDA代码中吗?
我想在CUDA C代码中使用汇编代码 为了减少昂贵的执行 就像我们在C编程中使用 asm 一样. 有可能吗? 解决方案 不,您不能,没有什么比C/C ++的ASM构造.您能做的就是调整生成的PTX组件,然后与CUDA一起使用. 请参阅 this 以示例. 但对于GPU,没有必要进行组装优化,您应该首先进行其他优化,例如记忆合并和占用率.请参阅其他解决方案 由于CUDA 4.0,CUDA工具链支持Inline PTX.工具包中有一个文档来描述它:using_inline_ptx_assembly_in_cuda.pdf 以下是一些代码,证明了在CUDA 4.0中使用内联PTX的使用.请注意,此代码不应用作CUDA内置__clz()函数的替代,我只是写了它来探索新的内联PTX功能的各个方面. __device__ __forceinline__ int my_clz (unsigned int x) { int res; asm ("
18 2023-02-04
编程技术问答社区
将PTX程序直接传递给CUDA驱动
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字符串)? 解决方案 取自ptxjit cuda示例: 将PTX程序定义为C字符串为 char myPtx32[] = "\n\ .version 1.4\n\ .ta
62 2023-02-03
编程技术问答社区
我怎样才能创建一个可执行文件来运行给定的PTX文件中的内核?
据我所知,您需要一个主机代码(用于CPU)和设备代码(对于GPU),没有它们,您无法在GPU上运行一些东西. 我正在学习PTX ISA,我不知道如何在Windows上执行它.我需要一个.cu文件来运行它,还是有另一种方式来运行它? 解决方案 tl; DR: 如何组装.ptx文件和主机代码文件并制作可执行文件? 你使用 cuda driver api .相关的示例代码是 vectorAddDrv (或许或许任何其他驱动程序API示例代码)以及 ptxjit . 我需要.cu文件来运行它,也需要另一种方式来运行它吗? 如果您在PTX表单中以设备代码开头,则不需要.cu文件(也不需要nvcc)来使用驱动程序API方法. 详细信息: 此答案的其余部分并非旨在成为驱动程序API编程的教程(使用已经给出的引用以及API参考手册此处),也不是ptx编程的教程.对于PTX编程,我将您引用 ptx文档. 要从开始,我们需要适当的PTX内核定义. (因为那
24 2023-02-01
编程技术问答社区
有可能把汇编指令放到CUDA代码中吗?
我想在CUDA C代码中使用汇编代码 为了减少昂贵的执行 正如我们在C编程中使用 asm 所做的那样. 是可能的吗? 解决方案 no,您不能,没有任何类似于C/C ++的ASM构造.你能做的是调整生成的PTX组件,然后用CUDA使用它. 看这个一个例子. 但对于GPU,没有必要的装配优化,您应该首先做其他优化,例如记忆聚合和占用.请参阅 cuda最佳实践指南有关更多信息. 其他解决方案 自CUDA 4.0以来,CUDA Toolchain支持内联PTX.工具包中有一个文档,它描述了它:使用_inline_ptx_assembly_in_cuda.pdf 以下是一些代码,演示在CUDA 4.0中使用内联PTX.请注意,此代码不应用作CUDA内置__CLZ()函数的替代品,我仅写它来探索新的内联PTX能力的方面. __device__ __forceinline__ int my_clz (unsigned int x) { int res;
34 2023-01-18
编程技术问答社区
漏斗转移|什么是漏斗转移?
通过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. ...但是它仍然无法解释"左漏斗移动"或"右漏斗变速"是什么. 那么,它是什么,一个人需要它? 解决方案 在CUDA的情况下,将两个32位寄存器串联为64位值.该值向左或向右移动;返回32位的最显着(
72 2022-12-01
编程技术问答社区
我是否应该研究一下PTX来优化我的内核?如果是这样,怎么做?
您是否建议阅读内核的PTX代码以进一步优化内核? 一个示例:我读到,如果自动循环展开工作,则可以从PTX代码中找出.如果不是这种情况,则必须在内核代码中手动展开循环. PTX代码还有其他用例吗? 您是否查看PTX代码? 我在哪里可以找到如何读取我的内核生成PTX代码的? 解决方案 关于PTX的第一个点是,它只是GPU上运行的代码的 InterMediate 表示 - 虚拟机汇编语言. PTX在编译时通过ptxas或在运行时由驱动程序组装到目标机器代码.因此,当您查看PTX时,您正在查看编译器发射的内容,但没有在GPU实际运行的内容上发出的内容.也可以编写您自己的PTX代码,要么从头开始(这是CUDA中支持的唯一jit汇编模型),要么是CUDA C代码中的Inline-Assembler部分的一部分(自CUDA 4.0以来已正式支持,但",但"非正式地"支持更长的时间". CUDA始终使用该工具包进行了完整的PTX语言指南,并已充分记录. 如果您想查看GPU实际运行
36 2022-10-23
编程技术问答社区
在CUDA 9中,一些本征被命名为"_sync()";语义相同吗?
在CUDA 9中,Nvidia似乎有了"合作团体"的新概念.由于某些原因,我对我来说并不完全清楚,__ballot()现在(= CUDA 9)对__ballot_sync()>的弃用.是别名还是语义更改了? ...对于现在已添加到其名称中的其他内置的类似问题. 解决方案 没有语义不是相同的.该功能本身是不同的,一个不是另一个的别名,新功能已经暴露出来,现在的实现行为在Volta架构和以前的体系结构之间有所不同. 首先,要设定地面工作,有必要认识到Volta 介绍了可能性 独立线程调度 .因此,Volta有可能在长时间内以非瓦普同步行为行为,并且在执行期间,以前的架构仍然可能是扭曲的同步. 仅通过为实际参与的线程提供预期的结果(即在该周期中的该指令问题上实际上是活跃的).现在,可以通过新的mask参数明确有关预期将参与哪些线程的程序员.但是,有一些要求,尤其是对Pascal和以前的体系结构.来自 请注意,对于Pascal和较早的架构,mask中的所有线程都必须执
16 2022-08-15
编程技术问答社区