内联 PTX 汇编代码强大吗?

2024-05-16

我看到一些代码示例,人们在 C 代码中使用内联 PTX 汇编代码。 CUDA工具包中的文档提到PTX很强大,为什么会这样呢?如果我们在 C 代码中使用这样的代码,我们会得到什么好处?


内联 PTX 使您可以访问未通过 CUDA 内在函数公开的指令,并允许您应用编译器中缺少或语言规范禁止的优化。有关使用内联 PTX 具有优势的有效示例,请参阅:cuda上的128位整数? https://stackoverflow.com/questions/6162140/128-bit-integer-on-cuda/6220499#6220499

使用内联 PTX 的 128 位加法仅需要 4 条指令,因为它可以直接访问进位标志。作为 HLL,C/C++ 没有进位标志的表示,因为给定的硬件平台可能没有进位标志(例如 MIPS)、单个进位标志(例如 x86、sm_2x),甚至多个进位标志。与 128 位加法和减法的 4 指令 PTX 版本相比,这些操作可以用 C 语言编码,如下所示:

#define SUBCcc(a,b,cy,t0,t1,t2) \
  (t0=(b)+cy, t1=(a), cy=t0<cy, t2=t1<t0, cy=cy+t2, t1-t0)
#define SUBcc(a,b,cy,t0,t1) \
  (t0=(b), t1=(a), cy=t1<t0, t1-t0)
#define SUBC(a,b,cy,t0,t1) \
  (t0=(b)+cy, t1=(a), t1-t0)
#define ADDCcc(a,b,cy,t0,t1) \
  (t0=(b)+cy, t1=(a), cy=t0<cy, t0=t0+t1, t1=t0<t1, cy=cy+t1, t0=t0)
#define ADDcc(a,b,cy,t0,t1) \
  (t0=(b), t1=(a), t0=t0+t1, cy=t0<t1, t0=t0)
#define ADDC(a,b,cy,t0,t1) \
  (t0=(b)+cy, t1=(a), t0+t1)

unsigned int cy, t0, t1, t2;

res.x = ADDcc  (augend.x, addend.x, cy, t0, t1);
res.y = ADDCcc (augend.y, addend.y, cy, t0, t1);
res.z = ADDCcc (augend.z, addend.z, cy, t0, t1);
res.w = ADDC   (augend.w, addend.w, cy, t0, t1);

res.x = SUBcc  (minuend.x, subtrahend.x, cy, t0, t1);
res.y = SUBCcc (minuend.y, subtrahend.y, cy, t0, t1, t2);
res.z = SUBCcc (minuend.z, subtrahend.z, cy, t0, t1, t2);
res.w = SUBC   (minuend.w, subtrahend.w, cy, t0, t1);

上面的加法和减法编译后的指令数可能是相应内联 PTX 版本使用的指令数的三到四倍。

本文内容由网友自发贡献,版权归原作者所有,本站不承担相应法律责任。如您发现有涉嫌抄袭侵权的内容,请联系:hwhale#tublm.com(使用前将#替换为@)

内联 PTX 汇编代码强大吗? 的相关文章

  • cuda中的count3非常慢

    我在 CUDA 中编写了一个小程序 用于计算 C 数组中有多少个 3 并打印它们 include
  • 某些子网格未使用 CUDA 动态并行执行

    我正在尝试 CUDA 5 0 GTK 110 中的新动态并行功能 我遇到了一个奇怪的行为 即我的程序没有返回某些配置的预期结果 不仅是意外的 而且每次启动都会出现不同的结果 现在我想我找到了问题的根源 似乎当生成太多子网格时 某些子网格 由
  • 使用 CUDA __device__ 函数时出现链接器错误 2005 和 1169(多重定义的符号)(默认情况下应内联)

    这个问题与以下问题有很大关系 A 如何将CUDA代码分成多个文件 https stackoverflow com questions 2090974 how to separate cuda code into multiple files
  • cudaMemcpyToSymbol 的问题

    我正在尝试复制到恒定内存 但我不能 因为我对 cudaMemcpyToSymbol 函数的用法有误解 我正在努力追随this http developer download nvidia com compute cuda 4 1 rel t
  • 为什么 gcc 和 NVCC (g++) 会看到两种不同的结构大小?

    我正在尝试将 CUDA 添加到 90 年代末编写的现有单线程 C 程序中 为此 我需要混合两种语言 C 和 C nvcc 是 c 编译器 问题在于 C 编译器将结构视为特定大小 而 C 编译器将相同的结构视为略有不同的大小 那很糟 我对此感
  • CUDA程序导致nvidia驱动程序崩溃

    当我超过大约 500 次试验和 256 个完整块时 我的 monte carlo pi 计算 CUDA 程序导致我的 nvidia 驱动程序崩溃 这似乎发生在 monteCarlo 内核函数中 任何帮助都会受到赞赏 include
  • cuda cpu功能-gpu内核重叠

    我在尝试开发以练习 CUDA 的 CUDA 应用程序时遇到并发问题 我想通过使用 cudaMemecpyAsync 和 CUDA 内核的异步行为来共享 GPU 和 CPU 之间的工作 但我无法成功重叠 CPU 执行和 GPU 执行 它与主机
  • 如何并行从数组中删除零值

    如何使用 CUDA 并行有效地从数组中删除零值 有关零值数量的信息是预先可用的 这应该可以简化这项任务 重要的是数字必须保持源数组中的顺序 当被复制到结果数组时 Example 该数组将例如包含以下值 0 0 19 7 0 3 5 0 0
  • CUDA Visual Studio 2010 Express 构建错误

    我正在尝试在 64 位 Windows 7 上使用 Visual Studio 2010 Express 在 Windows 上开始 CUDA 编程 我花了一段时间来设置环境 然后我刚刚编写了我的第一个程序 helloWorld cu 目前
  • CUDA 估计 2D 网格数据的每块线程数和块数

    首先我要说的是 我已经仔细阅读了所有类似的问题 确定每个块的线程和每个网格的块 https stackoverflow com questions 4391162 cuda determining threads per block blo
  • __syncthreads() 死锁

    如果只有部分线程执行 syncthreads 会导致死锁吗 我有一个这样的内核 global void Kernel int N int a if threadIdx x
  • 在 __device/global__ CUDA 内核中动态分配内存

    根据CUDA 编程指南 http developer download nvidia com compute cuda 3 2 prod toolkit docs CUDA C Programming Guide pdf 第 122 页 可
  • 尝试构建我的 CUDA 程序时出现错误 MSB4062

    当我尝试构建我的第一个 GPU 程序时 出现以下错误 有什么建议可能会出什么问题吗 错误 1 错误 MSB4062 Nvda Build CudaTasks SanitizePaths 任务 无法从程序集 C Program 加载 文件 M
  • cuda中有模板化的数学函数吗? [复制]

    这个问题在这里已经有答案了 我一直在寻找 cuda 中的模板化数学函数 但似乎找不到 在普通的 C 中 如果我调用std sqrt它是模板化的 并且将根据参数是浮点数还是双精度数执行不同的版本 我想要这样的 CUDA 设备代码 我的内核将真
  • cuda中内核的并行执行

    可以说我有三个全局数组 它们已使用 cudaMemcpy 复制到 GPU 中 但 c 中的这些全局数组尚未使用 cudaHostAlloc 分配 以便分配页面锁定的内存 而不是简单的全局分配 int a 100 b 100 c 100 cu
  • __device__ __constant__ 常量

    有什么区别吗 在 CUDA 程序中定义设备常量的最佳方法是什么 在 C 主机 设备程序中 如果我想将常量定义在设备常量内存中 我可以这样做 device constant float a 5 constant float a 5 问题 1
  • 将 nvidia 运行时添加到 docker 运行时

    我正在运行虚拟机GCP配备特斯拉 GPU 并尝试部署一个PyTorch基于应用程序使用 GPU 加速 我想让 docker 使用这个 GPU 可以从容器访问它 我设法在主机上安装了所有驱动程序 并且该应用程序在那里运行良好 但是当我尝试在
  • 通过 cuFFT 进行逆 FFT 缩放

    每当我使用 cuFFT 绘制程序获得的值并将结果与 Matlab 的结果进行比较时 我都会得到相同形状的图形 并且最大值和最小值位于相同的点 然而 cuFFT 得到的值比 Matlab 得到的值大得多 Matlab代码是 fs 1000 s
  • CUDA 中指令重放的其他原因

    这是我从 nvprof CUDA 5 5 获得的输出 Invocations Metric Name Metric Description Min Max Avg Device Tesla K40c 0 Kernel MyKernel do
  • 大型跨平台软件项目的技巧/资源

    我将开始一个大型软件项目 涉及跨平台 GUI 和大量的数字运算 我计划用 C 和 CUDA 编写大部分应用程序后端 并用 Qt4 编写 GUI 我计划使用 Make 作为我的构建系统 这将是一个只有两名开发人员的项目 一旦我相对深入地了解它

随机推荐