CUDA Cooperative Groups 例子

news2024/10/4 23:14:49

CUDA Cooperative Groups 例子

  • 一.复现步骤
  • 二.输出

CUDA Cooperative Groups是CUDA编程模型中引入的一组高级特性，旨在提供更灵活的线程组织和同步机制。通过Cooperative Groups，开发者可以在不同层次上组织线程，并执行更高效的并行操作。包括：

  • 网格组（Grid Group）：包含整个网格中所有线程的组。
  • 线程块组（Block Group）：包含线程块中所有线程的组。
  • 瓦片组（Tile Group）：将线程块划分为更小的线程子组，称为瓦片。

下文包含的测例:

  • 测试一:借助grid_group同步,将tid=0的数据复制给其它线程
  • 测试二:借助thread_block_tile同步,将每个thread block中的数据倒排
  • 测试三:tile内和
  • 测试四:tile内广播

一.复现步骤

tee cooperative_groups.cu<<-'EOF'
#include <iostream>
#include <cuda_runtime.h>
#include <iostream>
#include <stdio.h>
#include <assert.h>
#include <cstdio>
#include <cuda.h>

#include <cooperative_groups.h>
#include <cooperative_groups/reduce.h>

namespace cg = cooperative_groups;

#define CHECK_CUDA(call)                      \
  do {                              \
    cudaError_t err = call;                  \
    if (err != cudaSuccess) {                 \
      std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__; \
      std::cerr << " code=" << err << " (" << cudaGetErrorString(err) << ")" << std::endl; \
      exit(EXIT_FAILURE);                  \
    }                             \
  } while (0)

__device__ float gdata = 0;

/*
测试一:借助grid_group同步,将tid=0的数据复制给其它线程
*/
__global__ void case_0(float *iodata)
{
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  cg::grid_group grid = cg::this_grid();  
  if(tid==0) gdata=iodata[tid];
  grid.sync();
  iodata[tid]=gdata;
}

/*
测试二:借助thread_block_tile同步,将每个thread block中的数据倒排
*/
__global__ void case_1(float *iodata)
{
  __shared__ float sharedData[256];
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  cg::thread_block block = cg::this_thread_block();
  sharedData[threadIdx.x] = iodata[tid];
  block.sync();
  iodata[tid]=sharedData[blockDim.x-1-threadIdx.x];
}

/*
测试三:tile内和
*/
__global__ void case_2(float *iodata)
{
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  cg::thread_block block = cg::this_thread_block();
  cg::thread_block_tile<2> tile2 = cg::tiled_partition<2>(block);
  float sum = cg::reduce(tile2, iodata[tid], cg::plus<float>());
  tile2.sync();
  iodata[tid]=sum;
}

/*
测试三:tile内交换数据
*/
__global__ void case_3(float *iodata)
{
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  cg::thread_block block = cg::this_thread_block();
  cg::thread_block_tile<2> tile2 = cg::tiled_partition<2>(block);
  float nextValue = tile2.shfl(iodata[tid], (tile2.thread_rank() + 1) % tile2.size());
  tile2.sync();
  iodata[tid]=nextValue;
}

/*
测试四:tile内广播
*/
__global__ void case_4(float *iodata)
{
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  cg::thread_block block = cg::this_thread_block();
  cg::thread_block_tile<4> tile4 = cg::tiled_partition<4>(block);
  float value;
  //lane 1广播给其它lane
  if (tile4.thread_rank() == 1) {
     value = iodata[tid];
  }  
  value = tile4.shfl(value, 1);
  tile4.sync();
  iodata[tid]=value;
}

int main(int argc,char *argv[])
{
  int deviceid=0;cudaSetDevice(deviceid); 
  {
      printf(" ----------------- case 0 ----------------- \n");
      int block_count=4;
      int block_size=4;
      int thread_size=block_count*block_size;
      float *iodata;
      CHECK_CUDA(cudaHostAlloc(&iodata,thread_size*sizeof(float),cudaHostAllocDefault));  
      for(int i=0;i<thread_size;i++) iodata[i]=i+100;
      void *kernelArgs[] = {&iodata};
      cudaLaunchCooperativeKernel((void*)case_0, block_count, block_size, kernelArgs);
      CHECK_CUDA(cudaDeviceSynchronize());
      for(int i=0;i<thread_size;i++)
      {
        printf("tid:%02d %6.2f\n",i,iodata[i]);
      }
      CHECK_CUDA(cudaFreeHost(iodata));
  }
  {
      printf(" ----------------- case 1 ----------------- \n");
      int block_count=2;
      int block_size=4;
      int thread_size=block_count*block_size;
      float *iodata;
      CHECK_CUDA(cudaHostAlloc(&iodata,thread_size*sizeof(float),cudaHostAllocDefault));  
      for(int i=0;i<thread_size;i++) iodata[i]=i+100;
      void *kernelArgs[] = {&iodata};
      cudaLaunchCooperativeKernel((void*)case_1, block_count, block_size, kernelArgs);
      CHECK_CUDA(cudaDeviceSynchronize());
      for(int i=0;i<thread_size;i++)
      {
        printf("tid:%02d %6.2f\n",i,iodata[i]);
      }
      CHECK_CUDA(cudaFreeHost(iodata));
  }  
  {
      printf(" ----------------- case 2 ----------------- \n");
      int block_count=2;
      int block_size=8;
      int thread_size=block_count*block_size;
      float *iodata;
      CHECK_CUDA(cudaHostAlloc(&iodata,thread_size*sizeof(float),cudaHostAllocDefault));  
      for(int i=0;i<thread_size;i++) iodata[i]=i;
      void *kernelArgs[] = {&iodata};
      cudaLaunchCooperativeKernel((void*)case_2, block_count, block_size, kernelArgs);
      CHECK_CUDA(cudaDeviceSynchronize());
      for(int i=0;i<thread_size;i++)
      {
        printf("tid:%02d %6.2f\n",i,iodata[i]);
      }
      CHECK_CUDA(cudaFreeHost(iodata));
  }    
  {
      printf(" ----------------- case 3 ----------------- \n");
      int block_count=2;
      int block_size=8;
      int thread_size=block_count*block_size;
      float *iodata;
      CHECK_CUDA(cudaHostAlloc(&iodata,thread_size*sizeof(float),cudaHostAllocDefault));  
      for(int i=0;i<thread_size;i++) iodata[i]=i;
      void *kernelArgs[] = {&iodata};
      cudaLaunchCooperativeKernel((void*)case_3, block_count, block_size, kernelArgs);
      CHECK_CUDA(cudaDeviceSynchronize());
      for(int i=0;i<thread_size;i++)
      {
        printf("tid:%02d %6.2f\n",i,iodata[i]);
      }
      CHECK_CUDA(cudaFreeHost(iodata));
  }
  {
      printf(" ----------------- case 4 ----------------- \n");
      int block_count=2;
      int block_size=8;
      int thread_size=block_count*block_size;
      float *iodata;
      CHECK_CUDA(cudaHostAlloc(&iodata,thread_size*sizeof(float),cudaHostAllocDefault));  
      for(int i=0;i<thread_size;i++) iodata[i]=i;
      void *kernelArgs[] = {&iodata};
      cudaLaunchCooperativeKernel((void*)case_4, block_count, block_size, kernelArgs);
      CHECK_CUDA(cudaDeviceSynchronize());
      for(int i=0;i<thread_size;i++)
      {
        printf("tid:%02d %6.2f\n",i,iodata[i]);
      }
      CHECK_CUDA(cudaFreeHost(iodata));
  }  
}
EOF
/usr/local/cuda/bin/nvcc -std=c++17 -arch=sm_86 -lineinfo -o cooperative_groups cooperative_groups.cu \
 -I /usr/local/cuda/include -L /usr/local/cuda/lib64 -lcuda
./cooperative_groups

二.输出

 ----------------- case 0 -----------------
tid:00 100.00
tid:01 100.00
tid:02 100.00
tid:03 100.00
tid:04 100.00
tid:05 100.00
tid:06 100.00
tid:07 100.00
tid:08 100.00
tid:09 100.00
tid:10 100.00
tid:11 100.00
tid:12 100.00
tid:13 100.00
tid:14 100.00
tid:15 100.00
 ----------------- case 1 -----------------
tid:00 103.00
tid:01 102.00
tid:02 101.00
tid:03 100.00
tid:04 107.00
tid:05 106.00
tid:06 105.00
tid:07 104.00
 ----------------- case 2 -----------------
tid:00   1.00
tid:01   1.00
tid:02   5.00
tid:03   5.00
tid:04   9.00
tid:05   9.00
tid:06  13.00
tid:07  13.00
tid:08  17.00
tid:09  17.00
tid:10  21.00
tid:11  21.00
tid:12  25.00
tid:13  25.00
tid:14  29.00
tid:15  29.00
 ----------------- case 3 -----------------
tid:00   1.00
tid:01   0.00
tid:02   3.00
tid:03   2.00
tid:04   5.00
tid:05   4.00
tid:06   7.00
tid:07   6.00
tid:08   9.00
tid:09   8.00
tid:10  11.00
tid:11  10.00
tid:12  13.00
tid:13  12.00
tid:14  15.00
tid:15  14.00
 ----------------- case 4 -----------------
tid:00   1.00
tid:01   1.00
tid:02   1.00
tid:03   1.00
tid:04   5.00
tid:05   5.00
tid:06   5.00
tid:07   5.00
tid:08   9.00
tid:09   9.00
tid:10   9.00
tid:11   9.00
tid:12  13.00
tid:13  13.00
tid:14  13.00
tid:15  13.00

本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若转载，请注明出处：http://www.coloradmin.cn/o/2188881.html

如若内容造成侵权/违法违规/事实不符，请联系多彩编程网进行投诉反馈，一经查实，立即删除！

相关文章

20241004给荣品RD-RK3588-AHD开发板刷Rockchip原厂的Android12时永不休眠的步骤

20241004给荣品RD-RK3588-AHD开发板刷Rockchip原厂的Android12时永不休眠的步骤

20241004给荣品RD-RK3588-AHD开发板刷Rockchip原厂的Android12时永不休眠的步骤 2024/10/4 19:22 1、 Z:\rk3588s4_3588a12\device\rockchip\common\device.mk ifeq ($(strip $(BOARD_HAVE_BLUETOOTH_RTK)), true) include hardware/realtek/rtkbt/rtkbt.mk endif ifeq ($(str…
阅读更多...
YouTube音视频合并批处理基于 FFmpeg的

YouTube音视频合并批处理基于 FFmpeg的

专门针对YouTube高品质分享处理的&#xff0c;将音频和视频合并。 首先下载ffmpeg.exe网上随便下载。 echo off title YouTube 音视频合并 20241004 echo 作者&#xff1a;xiaoshen echo 网站&#xff1a;http://www.xiaoshen.cn/ echo. set /p audio请将【音频】文件拖拽到此…
阅读更多...
⌈ 传知代码 ⌋ 将一致性正则化用于弱监督学习

⌈ 传知代码 ⌋ 将一致性正则化用于弱监督学习

&#x1f49b;前情提要&#x1f49b; 本文是传知代码平台中的相关前沿知识与技术的分享~ 接下来我们即将进入一个全新的空间&#xff0c;对技术有一个全新的视角~ 本文所涉及所有资源均在传知代码平台可获取 以下的内容一定会让你对AI 赋能时代有一个颠覆性的认识哦&#x…
阅读更多...
什么是 NVIDIA 机密计算？( 上篇 )

什么是 NVIDIA 机密计算？( 上篇 )

什么是机密计算？ 文章目录 前言1. 机密计算定义2. 机密计算有何独特之处？3. 机密计算是如何得名的4. 机密计算的工作原理是什么？5. 缩小安全边界6. 机密计算的使用案例7. 机密计算如何发展8. 加速机密计算9. 机密计算的下一步前言 机密计算是一种在计算机处理器的受保护区域…
阅读更多...
全网最详细kubernetes中的资源

全网最详细kubernetes中的资源

1、资源管理介绍 在kubernetes中&#xff0c;所有的内容都抽象为资源&#xff0c;用户需要通过操作资源来管理kubernetes。 kubernetes的本质上就是一个集群系统&#xff0c;用户可以在集群中部署各种服务。 所谓的部署服务&#xff0c;其实就是在kubernetes集群中运行一个个的…
阅读更多...
csp-j模拟三补题报告

csp-j模拟三补题报告

前言 今天题难&#xff0c;排名没进前十 &#xff08;“关于二进制中一的个数的研究与规律”这篇文章正在写&#xff09; 第一题 三个&#xff08;three&#xff09; 我的代码&#xff08;AC&#xff09; #include<bits/stdc.h> #define ll long long using namespac…
阅读更多...
快停止这种使用U盘的行为！

快停止这种使用U盘的行为！

前言 现在各行各业的小伙伴基本上都需要用电脑来办公了&#xff0c;你敢说你不需要用电脑办公&#xff1f; 啊哈哈哈&#xff0c;用iPad或者手机办公的也算。 有些小伙伴可能经常996&#xff0c;甚至有时候都是007。有时候到了下班时间&#xff0c;工作还没做完&#xff0c;…
阅读更多...
Python技巧：如何处理未完成的函数

Python技巧：如何处理未完成的函数

一、问题的提出 写代码的时候&#xff0c;我们有时候会给某些未完成的函数预留一个空位&#xff0c;等以后有时间再写具体内容。通常&#xff0c;大家会用 pass 或者 ... &#xff08;省略号&#xff09;来占位。这种方法虽然能让代码暂时不报错&#xff0c;但可能在调试的时候…
阅读更多...
精准翻译神器：英汉互译软件的卓越表现

精准翻译神器：英汉互译软件的卓越表现

英文作为目前世界上使用最广的一种语言&#xff0c;是的很多先进的科学文献或者一些大厂产品的说明书都有英文的版本。为了方便我们的阅读和学习&#xff0c;现在有不少支持翻译英汉互译的工具&#xff0c;今天我们就一起来讨论一下吧。 1.福昕中英在线翻译 链接直达>>…
阅读更多...
二叉树的前序遍历——非递归版本

二叉树的前序遍历——非递归版本

1.题目解析 题目来源&#xff1a;144.二叉树的前序遍历——力扣 测试用例 2.算法原理 前序遍历&#xff1a; 按照根节点->左子树->右子树的顺序遍历二叉树 二叉树的前序遍历递归版本十分简单&#xff0c;但是如果树的深度很深会有栈溢出的风险&#xff0c;这里的非递归…
阅读更多...
【论文笔记】DKTNet: Dual-Key Transformer Network for small object detection

【论文笔记】DKTNet: Dual-Key Transformer Network for small object detection

【引用格式】&#xff1a;Xu S, Gu J, Hua Y, et al. Dktnet: dual-key transformer network for small object detection[J]. Neurocomputing, 2023, 525: 29-41. 【网址】&#xff1a;https://cczuyiliu.github.io/pdf/DKTNet%20Dual-Key%20Transformer%20Network%20for%20s…
阅读更多...
vue3实现打字机的效果,可以换行

vue3实现打字机的效果,可以换行

之前看了很多文章,效果是实现了,就是没有自动换行的效果,参考了文章写了一个,先上个效果图,卡顿是因为模仿了卡顿的效果,还是很丝滑的 目录 效果图:代码如下 效果图: ![请添加图片描述](https://i-blog.csdnimg.cn/direct/d8ef33d83dd3441a87d6d033d9e7cafa.gif 代码如下 原…
阅读更多...
jmeter学习（8）结果查看

jmeter学习（8）结果查看

1&#xff09;查看结果树 查看结果树&#xff0c;显示取样器请求和响应的细节以及请求结果&#xff0c;包括消息头&#xff0c;请求的数据&#xff0c;响应的数据。 2&#xff09;汇总报告 汇总报告&#xff0c;为测试中的每个不同命名的请求创建一个表行。这与聚合报告类似&…
阅读更多...
[数据结构] 树

[数据结构] 树

n个结点的有限集合 除了根节点以外&#xff0c;每一个结点有且只有一条与父节点的连线&#xff1b; 总共有N-1条连线。 子树之间不相交。 术语 树的表示 每个结点的结构不知道 可以统一设置结构&#xff0c;优点&#xff1a;处理方便 缺点&#xff1a;会造成空间浪费&…
阅读更多...
Chromium 硬件加速开关c++

Chromium 硬件加速开关c++

选项页控制硬件加速开关 1、前端代码 <settings-toggle-button id"hardwareAcceleration"pref"{{prefs.hardware_acceleration_mode.enabled}}"label"$i18n{hardwareAccelerationLabel}"><template is"dom-if" if"[…
阅读更多...
6.5 监控和日志 架构模式和应用实践

6.5 监控和日志 架构模式和应用实践

6.5 监控和日志 架构模式和应用实践 目录概述需求&#xff1a; 设计思路实现思路分析1.集中式监控2.分布式监控3.边缘监控4.集中式日志管理5.分布式日志管理6.实时日志流处理 监控工具最佳实践 参考资料和推荐阅读 Survive by day and develop by night. talk for import biz ,…
阅读更多...
基于元神操作系统实现NTFS文件操作（六）

基于元神操作系统实现NTFS文件操作（六）

1. 背景 本文主要介绍$Root元文件属性的解析。先介绍元文件各属性的属性体构成&#xff0c;然后结合读取到的元文件内容&#xff0c;对测试磁盘中目标分区的根目录进行展示。 2. $Root元文件属性的解析 使用每个属性头偏移0x04-0x07处的值可以从第一个属性开始依次定位下一个…
阅读更多...
Jupyter | jupyter notebook 使用 conda 环境

Jupyter | jupyter notebook 使用 conda 环境

博客使用更佳 点我进入博客 创建虚拟环境 在 Anaconda Prompt 里面输入&#xff1a; conda create -n env-name并且输入 y 确认。例如我们创建环境名为 jupyter 激活环境 conda activate env-name激活之后发现环境从 base 变为 jupyter(笔者用的 env-name 为 jupyter) …
阅读更多...
python-求一个整数的质因数/字符串的镜像/加数

python-求一个整数的质因数/字符串的镜像/加数

一:求一个整数的质因数 题目描述 编写一个程序&#xff0c;返回给定整数的质因数。 定义函数get_prime_factors()&#xff0c;该函数接受一个参数num&#xff08;正整数&#xff09;。 该函数应返回传入参数的质因数列表&#xff0c;且从小到大排序。 比如150的质因数分解如…
阅读更多...
Spring MVC__HttpMessageConverter、拦截器、异常处理器、注解配置SpringMVC、SpringMVC执行流程

Spring MVC__HttpMessageConverter、拦截器、异常处理器、注解配置SpringMVC、SpringMVC执行流程

目录 一、HttpMessageConverter1、RequestBody2、RequestEntity3、ResponseBody4、SpringMVC处理json5、SpringMVC处理ajax6、RestController注解7、ResponseEntity7.1、文件下载7.2、文件上传 二、拦截器1、拦截器的配置2、拦截器的三个抽象方法3、多个拦截器的执行顺序 三、异…
阅读更多...
推荐文章
最新文章

Copyright @ 2022~2023