程序员经验分享-hwhale

Cuda—— Atomic Functions

2023-05-16

Atomic Functions

原子函数对驻留在全局或共享内存中的一个 32 位或 64 位字执行读取-修改-写入原子操作。 在 float2 或 float4 的情况下，对驻留在全局内存中的向量的每个元素执行读取-修改-写入操作。 例如，atomicAdd() 在全局或共享内存中的某个地址读取一个字，向其添加一个数字，然后将结果写回同一地址。 原子函数只能在设备函数中使用。

本节中描述的原子函数具有排序 cuda::memory_order_relaxed 并且仅在特定范围内是原子的：

  • 具有 _system 后缀的原子 API（示例：__atomicAdd_system）在 cuda::thread_scope_system 范围内是原子的。

  • 没有后缀的原子 API（例如：__atomicAdd）在 cuda::thread_scope_device 范围内是原子的。

  • 带有 _block 后缀的原子 API（例如：__atomicAdd_block）在 cuda::thread_scope_block 范围内是原子的。

在以下示例中，CPU 和 GPU 都自动更新地址 addr 处的整数值：

__global__ void mykernel(int *addr) {
  atomicAdd_system(addr, 10);       // only available on devices with compute capability 6.x
}

void foo() {
  int *addr;
  cudaMallocManaged(&addr, 4);
  *addr = 0;

   mykernel<<<...>>>(addr);
   __sync_fetch_and_add(addr, 10);  // CPU atomic operation
}

注意，任何原子操作都可以基于 atomicCAS() (Compare And Swap) 来实现。 例如，双精度浮点数的 atomicAdd() 在计算能力低于 6.0 的设备上不可用，但可以按如下方式实现：
查看计算能力：https://developer.nvidia.com/zh-cn/cuda-gpus#compute

#if __CUDA_ARCH__ < 600
__device__ double atomicAdd(double* address, double val)
{
    unsigned long long int* address_as_ull =
                              (unsigned long long int*)address;
    unsigned long long int old = *address_as_ull, assumed;

    do {
        assumed = old;
        old = atomicCAS(address_as_ull, assumed,
                        __double_as_longlong(val +
                               __longlong_as_double(assumed)));

    // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN)
    } while (assumed != old);

    return __longlong_as_double(old);
}
#endif

以下设备范围的原子 API 有系统范围和块范围的变体，但以下情况除外：

  • 计算能力小于 6.0 的设备只支持设备范围的原子操作，

  • 计算能力低于 7.2 的 Tegra 设备不支持系统范围的原子操作。

算术函数

atomicAdd()

int atomicAdd(int* address, int val);
unsigned int atomicAdd(unsigned int* address, unsigned int val);
unsigned long long int atomicAdd(unsigned long long int* address, unsigned long long int val);

读取位于全局或共享内存中address的 16 位、32 位或 64 位 old，计算 (old + val)，并将结果存储回同一地址的内存。 这三个操作在一个原子事务中执行。 该函数返回旧的。

float atomicAdd(float* address, float val);
double atomicAdd(double* address, double val);

atomicAdd() 的 32 位浮点版本仅受计算能力 2.x 及更高版本的设备支持。
atomicAdd() 的 64 位浮点版本仅受计算能力 6.x 及更高版本的设备支持。

__half2 atomicAdd(__half2 *address, __half2 val);
__half atomicAdd(__half *address, __half val);
__nv_bfloat162 atomicAdd(__nv_bfloat162 *address, __nv_bfloat162 val);
__nv_bfloat16 atomicAdd(__nv_bfloat16 *address, __nv_bfloat16 val);

atomicAdd() 的 32 位 __half2 浮点版本仅受计算能力 6.x 及更高版本的设备支持。 对于两个 __half 或 __nv_bfloat16 元素中的每一个，分别保证 __half2 或 __nv_bfloat162 添加操作的原子性； 作为单个 32 位访问，不能保证整个 __half2 或 __nv_bfloat162 是原子的。

atomicAdd() 的 16 位 __half 浮点版本仅受计算能力 7.x 及更高版本的设备支持。

atomicAdd() 的 16 位 __nv_bfloat16 浮点版本仅受计算能力 8.x 及更高版本的设备支持。

float2 atomicAdd(float2* address, float2 val);
float4 atomicAdd(float4* address, float4 val);

atomicAdd() 的 float2 和 float4 浮点向量版本仅受计算能力 9.x 及更高版本的设备支持。 float2 或 float4 add 操作的原子性分别为两个或四个 float 元素中的每一个保证； 作为单个 64 位或 128 位访问，不能保证整个 float2 或 float4 是原子的。

atomicAdd() 的 float2 和 float4 浮点向量版本仅受计算能力 9.x 及更高版本的设备支持。
atomicAdd() 的 float2 和 float4 浮点向量版本仅支持全局内存地址。

atomicSub()

int atomicSub(int* address, int val);
unsigned int atomicSub(unsigned int* address, unsigned int val);

读取位于全局或共享内存中地址地址的 32 位字 old，计算 (old - val)，并将结果存储回同一地址的内存。 这三个操作在一个原子事务中执行。 该函数返回旧的。

atomicExch()

int atomicExch(int* address, int val);
unsigned int atomicExch(unsigned int* address,　unsigned int val);
unsigned long long int atomicExch(unsigned long long int* address, unsigned long long int val);
float atomicExch(float* address, float val);

读取位于全局或共享内存地址处的 32 位或 64 位字 old，并将 val 存储回同一地址的内存。 这两个操作在一个原子事务中执行。 该函数返回旧的。

atomicMin()

int atomicMin(int* address, int val);
unsigned int atomicMin(unsigned int* address, unsigned int val);
unsigned long long int atomicMin(unsigned long long int* address, unsigned long long int val);
long long int atomicMin(long long int* address, long long int val);

读取位于全局或共享内存地址的 32 位或 64 位字 old，计算 old 和 val 的最小值，并将结果存储回同一地址的内存。 这三个操作在一个原子事务中执行。 该函数返回旧的。

64 位版本的 atomicMin() 仅受计算能力 5.0 及更高版本的设备支持。

atomicMax()

int atomicMax(int* address, int val);
unsigned int atomicMax(unsigned int* address, unsigned int val);
unsigned long long int atomicMax(unsigned long long int* address, unsigned long long int val);
long long int atomicMax(long long int* address, long long int val);

读取位于全局或共享内存地址处的 32 位或 64 位字 old，计算 old 和 val 的最大值，并将结果存储回同一地址的内存。 这三个操作在一个原子事务中执行。 该函数返回旧的。

64 位版本的 atomicMax() 仅受计算能力 5.0 及更高版本的设备支持。

atomicInc()

unsigned int atomicInc(unsigned int* address, unsigned int val);

读取位于全局或共享内存地址处的 32 位字 old，计算 ((old >= val) ? 0 : (old+1))，并将结果存储回同一地址的内存。 这三个操作在一个原子事务中执行。 该函数返回旧的。

atomicDec()

unsigned int atomicDec(unsigned int* address, unsigned int val);

读取位于全局或共享内存中地址地址的 32 位字 old，计算 (((old == 0) || (old > val)) ? val : (old-1) )，并将结果存回 到同一地址的内存。 这三个操作在一个原子事务中执行。 该函数返回旧的。

atomicCAS()

int atomicCAS(int* address, int compare, int val);
unsigned int atomicCAS(unsigned int* address,
                       unsigned int compare,
                       unsigned int val);
unsigned long long int atomicCAS(unsigned long long int* address,
                                 unsigned long long int compare,
                                 unsigned long long int val);
unsigned short int atomicCAS(unsigned short int *address,
                             unsigned short int compare,
                             unsigned short int val);

读取位于全局或共享内存中地址地址的 16 位、32 位或 64 位字 old，计算 (old == compare ? val : old)，并将结果存储回同一地址的内存。 这三个操作在一个原子事务中执行。 该函数返回旧的（比较和交换）。

Bitwise Functions

atomicAnd()

int atomicAnd(int* address, int val);
unsigned int atomicAnd(unsigned int* address,
                       unsigned int val);
unsigned long long int atomicAnd(unsigned long long int* address,
                                 unsigned long long int val);

读取位于全局或共享内存中地址address 的32 位或64 位字old，计算(old & val)，并将结果存储回同一地址的内存。 这三个操作在一个原子事务中执行。 该函数返回旧的。

atomicAnd() 的 64 位版本仅受计算能力 5.0 及更高版本的设备支持。

atomicOr()

int atomicOr(int* address, int val);
unsigned int atomicOr(unsigned int* address,
                      unsigned int val);
unsigned long long int atomicOr(unsigned long long int* address,
                                unsigned long long int val);

读取位于全局或共享内存中地址地址的 32 位或 64 位字 old，计算 (old | val)，并将结果存储回同一地址的内存。 这三个操作在一个原子事务中执行。 该函数返回旧的。

atomicOr() 的 64 位版本仅受计算能力 5.0 及更高版本的设备支持。

atomicXor()

int atomicXor(int* address, int val);
unsigned int atomicXor(unsigned int* address,
                       unsigned int val);
unsigned long long int atomicXor(unsigned long long int* address,
                                 unsigned long long int val);

读取位于全局或共享内存中地址地址的 32 位或 64 位字 old，计算 (old ^ val)，并将结果存储回同一地址的内存。 这三个操作在一个原子事务中执行。 该函数返回旧的。

atomicXor() 的 64 位版本仅受计算能力 5.0 及更高版本的设备支持。

参考：https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#bitwise-functions

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

CUDA

atomic

Functions

Cuda—— Atomic Functions 的相关文章

  • CUDA 矩阵加法时序，按行与按行比较按栏目

    我目前正在学习 CUDA 并正在做一些练习 其中之一是实现以 3 种不同方式添加矩阵的内核 每个元素 1 个线程 每行 1 个线程和每列 1 个线程 矩阵是方阵 并被实现为一维向量 我只需用以下命令对其进行索引 A N row col 直觉

  • 领域驱动设计：处理原子操作和事务

    必须保证每个聚合内部的一致性 在存储库中执行此操作很容易 因为我始终可以使用数据库或框架中的事务 我对存储库之外发生的事情表示怀疑 一项服务可能需要使用多个聚合来处理请求 在服务处理过程中或在保留聚合时可能会出现问题 如果服务处理过程中出现

  • 如何使用 CUDA/Thrust 对两个数组/向量根据其中一个数组中的值进行排序

    这是一个关于编程的概念问题 总而言之 我有两个数组 向量 我需要对一个数组 向量进行排序 并将更改传播到另一个数组 向量中 这样 如果我对 arrayOne 进行排序 则对于排序中的每个交换 arrayTwo 也会发生同样的情况 现在 我知

  • “gld/st_throughput”和“dram_read/write_throughput”指标之间有什么区别？

    在 CUDA 可视化分析器版本 5 中 我知道 gld st requested throughput 是应用程序请求的内存吞吐量 然而 当我试图找到硬件的实际吞吐量时 我很困惑 因为有两对似乎合格的指标 它们是 gld st throug

  • cuda-gdb 错误消息

    我尝试使用 cuda gdb 调试我的 CUDA 应用程序 但遇到了一些奇怪的错误 我设置了选项 g G O0构建我的应用程序 我可以在没有 cuda gdb 的情况下运行我的程序 但没有得到正确的结果 因此我决定使用 cuda gdb 但

  • 具有 Cuda Thrust 的多个 GPU？

    如何将 Thrust 与多个 GPU 一起使用 这只是使用 cudaSetDevice deviceId 的问题吗 然后运行相关的 Thrust 代码 使用 CUDA 4 0 或更高版本 cudaSetDevice deviceId 接下来

  • 有条件减少 CUDA

    我需要总结一下100000值存储在数组中 但带有条件 有没有办法在 CUDA 中做到这一点以快速产生结果 任何人都可以发布一个小代码来做到这一点吗 我认为 要执行条件约简 您可以直接将条件引入为乘法0 假 或1 真 加数 换句话说 假设您希

  • cudaDeviceScheduleBlockingSync 和 cudaDeviceScheduleYield 之间有什么区别？

    正如这里所说 如何减少 CUDA 同步延迟 延迟 https stackoverflow com questions 11953722 how to reduce cuda synchronize latency delay 等待设备结果有

  • Bank 在字长方面存在冲突

    我读过一些关于共享内存的好文章 但我对银行冲突有初步疑问 据说 如果线程 1 和线程 2 从存储体 0 访问字 0 则不存在存储体冲突 但如果他们访问不同的单词 就会出现银行冲突 但我的问题是不同的单词如何可以驻留在一个银行中 由于bank

  • 设备内存刷新cuda

    我正在运行一个 C 程序 其中调用了两次 cuda 主机函数 我想清理这两个调用之间的设备内存 有没有办法可以刷新 GPU 设备内存 我使用的是计算能力为2 0的Tesla M2050 如果你只想将内存归零 那么cudaMemset可能是最

  • OpenCV 2.4.3rc 和 CUDA 4.2：“OpenCV 错误：没有 GPU 支持”

    我在这张专辑中上传了几张截图 https i stack imgur com TELST jpg https i stack imgur com TELST jpg 我正在尝试在 Visual Studio 2008 中的 OpenCV 中

  • CUDA 的嵌套循环

    我想将我的 C 代码移植到 CUDA 主要计算部分包含3个for嵌套循环 for int i 0 i lt Nx i for int j 0 j

  • CUDA 模型 - 什么是扭曲尺寸？

    最大工作组大小和扭曲大小之间有什么关系 假设我的设备有 240 个 CUDA 流处理器 SP 并返回以下信息 CL DEVICE MAX COMPUTE UNITS 30 CL DEVICE MAX WORK ITEM SIZES 512

  • 使用 GPU 进行 Matlab 卷积

    我用gpuArray尝试了matlab的卷积函数conv2 convn 例如 convn gpuArray rand 100 100 10 single gpuArray rand 5 single 并将其与 cpu 版本 convn ra

  • 竞争条件和 Clojure Atoms

    clojure atom 的文档指出 Changes to atoms are always free of race conditions 然而 竞争条件不仅是根据更改定义的 而且是在不同线程中并行逻辑操作的上下文中定义的 我想知道 保证

  • 在 CUDA 中的设备内存上分配 2D 数组

    如何在 Cuda 中的设备内存中分配和传输 往返于主机 2D 数组 我找到了解决这个问题的方法 我不必展平阵列 内置的cudaMallocPitch 函数完成了这项工作 我可以使用以下命令将阵列传输到设备或从设备传输阵列cudaMemcpy

  • 我们如何在每次运行时使用不同的种子在 CUDA C 中生成随机数？

    我正在研究一个随机过程 我想在每次运行程序时在 CUDA 内核中生成不同的系列随机数 这类似于我们在 C 中声明 种子 时间 空 接下来是 srand 种子 和兰特 我可以通过内核将种子从主机传递到设备 但是 这样做的问题是我必须将整个种子

  • 如何安装libcusolver.so.11

    我正在尝试安装 Tensorflow 但它要求 libcusolver so 11 而我只有 libcusolver so 10 有人可以告诉我我做错了什么吗 这是我的 Ubuntu nvidia 和 CUDA 版本 uname a Lin

  • 如何在没有 nvcc 的情况下在编译时获取 CUDA 工具包版本？

    我在 cpp 文件中对 cuSPARSE 库进行了一些调用 这些调用在旧工具包中不可用 为了支持使用旧工具包的系统 我想使用编译器指令编译不同的代码部分 特别是 我想使用旧工具包的 CSR 格式矩阵和新工具包的 BSR 格式矩阵来求解稀疏三

  • CUDA-Kernel 应该根据块大小动态崩溃

    我想做稀疏矩阵 密集向量乘法 假设用于压缩矩阵中条目的唯一存储格式是压缩行存储 CRS 我的内核如下所示 global void krnlSpMVmul1 float data mat int num nonzeroes unsigned

随机推荐

  • CUDA动态并行实现快速排序

    简介 排序是任何应用的基本构造块的关键算法之一 有许多排序算法已经被广泛研究 xff0c 常见的排序算法时间和空间复杂度如下 xff1a 一些排序算法属于分治算法的范畴 这些算法适用于并行性 xff0c 并适合 GPU 等架构 xff0c

  • 3d稀疏卷积——spconv源码剖析（一）

    本节主要是介绍下卷积的理论基础 结合spconv代码剖析从第二小节开始介绍 xff0c 本节介绍2D和3D卷积基础理论和稀疏卷积分类 xff0c 后再详细介绍下3d稀疏卷积的工作原理 2D卷积 2D卷积 xff1a 卷积核在输入图像的二维空

  • 基于Spring Cloud Zuul搭建网关服务

    1 网关服务所谓何 在微服务架构风格中 xff0c 一个大应用被拆分成为了多个小的服务系统提供出来 xff0c 这些小的系统他们可以自成体系 xff0c 也就是说这些小系统可以拥有自己的数据库 xff0c 框架甚至语言等 xff0c 这些小

  • Redis 命令

    命令 描述Redis GEOADD 命令 将指定的地理空间位置 xff08 纬度 经度 名称 xff09 添加到指定的key中Redis GEODIST 命令 返回两个给定位置之间的距离Redis GEOHASH 命令 返回一个或多个位置元

  • 3d稀疏卷积——spconv源码剖析（二）

    本文基于OpenPCDet框架中CeneterPoint算法 xff0c 对spconv库中稀疏卷积源码进行剖析 xff1a 首先看OpenPCDet下的pcdet models backbones 3d spconv backbone p

  • 3d稀疏卷积——spconv源码剖析（三）

    构建Rulebook 下面看ops get indice pairs xff0c 位于 xff1a spconv ops py 构建Rulebook由ops get indice pairs接口完成 get indice pairs函数具体

  • 3d稀疏卷积——spconv源码剖析（四）

    普通3d稀疏卷积RuleBook构建 我们继续看普通稀疏卷积RuleBook的建立过程 xff0c 返回src spconv spconv ops cc 看getIndicePairs函数的普通3D稀疏卷积部分 span class tok

  • 3d稀疏卷积——spconv源码剖析（五）

    下面介绍了根据构建的Rulebook执行具体稀疏卷积计算 xff0c 继续看类SparseConvolution 代码位于 xff1a spconv conv py span class token keyword class span s

  • TensoRT API自定义trt网络结构

    这个后续有时间进一步整理 pth转wts 若使用tensorrt加载wts格式 xff0c 需将模型训练的pt pth ckpt等格式权重转换为wts span class token keyword def span span class

  • 生成voxel——spconv源码剖析（六）

    CPU 先看spconv1 0 中cpu版本的generate voxels xff0c spconv1 0无gpu版本 看centerpoint的预处理pcdet datasets processor data processor py

  • CUDA——向量化内存

    许多 CUDA 内核受带宽限制 xff0c 新硬件中触发器与带宽的比率增加导致更多带宽受限内核 这使得采取措施缓解代码中的带宽瓶颈变得非常重要 在本文中 xff0c 我将向您展示如何在 CUDA C C 43 43 中使用矢量加载和存储来帮

  • TensoRT—— buffers管理（samplesCommon::BufferManager）

    BufferManager类处理主机和设备buffer分配和释放 这个RAII类处理主机和设备buffer的分配和释放 主机和设备buffers之间的memcpy以帮助inference xff0c 以及debugging dumps以验证

  • cuda Sgemm矩阵乘法优化

    近期在部署3d稀疏卷积 xff0c 需要对Rulebook与weights的计算进行速度优化 xff0c 先研究下cuda矩阵乘法 xff0c 特此记录下 xff1a CPU span class token keyword void sp

  • 基于pointpillars的点云目标检测、测试评估、TensorRT后量化及ROS可视化

    代码已经开源 xff1a https github com Xiao Hu Z pointpillars int8 安装环境 Prepare the OpenPCDet environment 导出onnx To export your o

  • 用pip安装pymongo模块报错：Could not find a version that satisfies the requirement pymongo(from version:)

    cmd中想用pip安装pymongo模块 显示黄色错误信息ReadTimeoutError和ConnectTimeoutError Retrying几次之后显示红色错误信息Could not find a version that sati

  • TensorRT量化工具pytorch_quantization代码解析(一）

    量化工具箱pytorch quantization 通过提供一个方便的 PyTorch 库来补充 TensorRT xff0c 该库有助于生成可优化的 QAT 模型 该工具包提供了一个 API 来自动或手动为 QAT 或 PTQ 准备模型

  • TensorRT量化工具pytorch_quantization代码解析(二）

    后续继续补充 xff01 继续看张量量化函数 xff0c 代码位于 xff1a tools pytorch quantization pytorch quantization tensor quant py ScaledQuantDescr

  • TensorRT量化工具pytorch_quantization代码解析(四）

    继续看pytorch quantiation calib 中Calibrator类 xff0c 代码位于 xff1a tools pytorch quantization pytorch quantization calib 其作用 xff

  • 基于3d稀疏卷积的centerpoint部署

    目前已实现基于稀疏卷积的centerpoint部署 xff0c 精度不丢失 xff0c 在3080ti 显卡nuscenes数据集下 pytorch 平均一帧耗时182ms xff0c 本文部署centerpoint fp32推理平均一帧耗

  • Cuda—— Atomic Functions

    Atomic Functions 原子函数对驻留在全局或共享内存中的一个 32 位或 64 位字执行读取 修改 写入原子操作 在 float2 或 float4 的情况下 xff0c 对驻留在全局内存中的向量的每个元素执行读取 修改 写入操

热门标签