CUDA Cooperative Groups 例子

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广播给其它laneif (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.xdnf.cn/news/1551077.html

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

相关文章

Golang | Leetcode Golang题解之第447题回旋镖的数量

Golang | Leetcode Golang题解之第447题回旋镖的数量

题目&#xff1a; 题解&#xff1a; func numberOfBoomerangs(points [][]int) (ans int) {for _, p : range points {cnt : map[int]int{}for _, q : range points {dis : (p[0]-q[0])*(p[0]-q[0]) (p[1]-q[1])*(p[1]-q[1])cnt[dis]}for _, m : range cnt {ans m * (m - 1)…
阅读更多...
Vue3 + Vite 开发环境下解决跨域问题：配置代理服务器

Vue3 + Vite 开发环境下解决跨域问题：配置代理服务器

一、介绍 在 Vue3 结合 Vite 的前端开发中&#xff0c;跨域问题是常见的挑战之一。特别是在开发阶段&#xff0c;当后端 API 尚未配置好 CORS 支持时&#xff0c;使用代理服务器来绕过浏览器的同源策略&#xff08;Same-origin policy&#xff09;就显得尤为重要。本文将介绍如…
阅读更多...
Word办公自动化的一些方法

Word办公自动化的一些方法

1.Word部分内容介绍 word本身是带有格式的一种文档&#xff0c;有人说它本质是XML&#xff0c;所以一定要充分利用标记了【样式】的特性来迅速调整【格式】&#xff0c;从而专心编辑文档内容本身。 样式&#xff08;集&#xff09; 编号&#xff08;多级关联样式编号&#xff…
阅读更多...
Ubuntu下安装Zookeeper集群

Ubuntu下安装Zookeeper集群

Zookeeper集群是一个开源的分布式协调服务系统&#xff0c;它由Apache软件基金会维护&#xff0c;旨在为分布式应用提供一致性和可靠性的服务。 在Zookeeper集群中&#xff0c;服务器可以扮演三种角色——领导者&#xff08;Leader&#xff09;、跟随者&#xff08;Follower&a…
阅读更多...
去掉顶部下拉框中的护眼模式

去掉顶部下拉框中的护眼模式

文章目录 需求参考知识点功能实现问题点问题点修改文件具体修改点总结 需求 顶部下拉框中的护眼模式去掉 参考知识点 我们在博客中已经分析了QS、QQS 相关的流程&#xff0c;实践过项目需求 Android12 SystemUI QS面板新增截屏功能 Android12_SystemUI下拉框新增音量控制条…
阅读更多...
微信网页 上传图片压缩

微信网页 上传图片压缩

微信网页上传图片时的压缩问题可以通过多种方法解决。以下是一些有效的方案和相关API的使用说明。 主要解决方案 1. 使用Canvas进行自定义压缩: 对于需要适配多种设备和格式的情况，可以利用Canvas API进行图片重绘和压缩。通过获取图片信息、设置Canvas尺寸、绘制图片并…
阅读更多...
2024年9月30日历史上的今天大事件早读

2024年9月30日历史上的今天大事件早读

1626年9月30日 清太祖努尔哈赤去世 1862年9月30日 德国首任宰相俾斯麦实行“铁血政策” 1887年9月30日 黄河决口 1931年9月30日 国际联盟决议日本撤兵 1937年9月30日 平型关战役结束 1938年9月30日 慕尼黑协议签订 1938年9月30日 前中华民国国务总理唐绍仪遇刺身亡 1941…
阅读更多...
box-im

box-im

任何一个开源项目&#xff0c;都可以让自己得到提升&#xff01; 启动minio: minio.exe server C:\Program Files\Minio\minioData
阅读更多...
相亲交友系统源码中的数据安全策略

相亲交友系统源码中的数据安全策略

在数字化时代&#xff0c;人们越来越依赖于互联网来寻找生活的另一半。相亲交友系统作为连接单身男女的重要平台&#xff0c;承载着无数用户的个人信息与隐私。因此&#xff0c;数据安全成为了此类系统不可忽视的关键因素。本文将探讨相亲交友系统源码中的数据安全策略&#xf…
阅读更多...
事实与价值双阈值是算计启动的门槛

事实与价值双阈值是算计启动的门槛

在现代社会&#xff0c;个体与群体的决策过程受到多种因素的影响&#xff0c;其中事实与价值的关系尤为重要。事实作为客观存在的基础&#xff0c;价值则是主观认知的体现。两者的相互作用构成了人类行为的复杂性&#xff0c;尤其在经济学、社会学以及伦理学等领域&#xff0c;…
阅读更多...
3GPP链路级仿真-Link-Level Simulator for 5G Localization

3GPP链路级仿真-Link-Level Simulator for 5G Localization

文章目录 II. SYSTEM ARCHITECTURE AND CAPABILITIESA. System Architecture III. KEY COMPONENTSA. Transmission Models of the Positioning SignalsB. Dedicated Wireless Channel Model IV. APPLICATION CASESA. Two-Dimensional Mobile Terminal Localization仿真工作流程…
阅读更多...
合成孔径雷达海上石油泄露分割数据集，共8000对图像，sentinel和palsar传感器，共400MB

合成孔径雷达海上石油泄露分割数据集，共8000对图像，sentinel和palsar传感器，共400MB

合成孔径雷达海上石油泄露分割数据集&#xff0c;共8000对图像&#xff0c;sentinel和palsar传感器&#xff0c;共400MB 名称 合成孔径雷达&#xff08;SAR&#xff09;海上石油泄露分割数据集 规模 图像对数&#xff1a;8000对图像传感器类型&#xff1a; Sentinel-1 SAR 传…
阅读更多...
PTVS：Python开发者的Visual Studio扩展

PTVS：Python开发者的Visual Studio扩展

PTVS&#xff08;Python Tools for Visual Studio&#xff09; 是由微软开发的一款用于 Visual Studio 的 Python 开发环境。它为 Python 开发者提供了丰富的功能&#xff0c;包括智能感知、调试工具、项目管理、测试支持等&#xff0c;使得在 Windows 平台上进行 Python 开发变…
阅读更多...
光耦——为智能电网建设提供安全高效解决方案

光耦——为智能电网建设提供安全高效解决方案

在智能电网的宏伟蓝图中&#xff0c;光耦以其独特的光电转换特性&#xff0c;成为了不可或缺的技术核心。作为一种先进的光电转换器件&#xff0c;光耦在智能电网建设中发挥着举足轻重的作用。 电能计量与监测 光耦可用于智能电表等电能计量设备中&#xff0c;实现电能数据的采…
阅读更多...
数据库管理-第245期 主流国产数据库RAC架构概览（20240929）

数据库管理-第245期 主流国产数据库RAC架构概览（20240929）

数据库管理245期 2024-09-29 数据库管理-第245期 主流国产数据库RAC架构概览&#xff08;20240929&#xff09;1 DMDSC2 KingBaseES RAC3 PolarDB4 Cantian5 HaloDB DLB/Data Sharding总结 数据库管理-第245期 主流国产数据库RAC架构概览&#xff08;20240929&#xff09; 作者…
阅读更多...
流量卡领取完就下架怎么回事，这种情况还能用吗？

流量卡领取完就下架怎么回事，这种情况还能用吗？

流量卡领取完就下架怎么回事&#xff0c;这种情况还能用吗&#xff1f;可以&#xff0c;这是很正常的一种情况&#xff01; 大家可能都发现了&#xff0c;网上的流量卡虽然月租低、流量多&#xff0c;但是却有一个弊端&#xff0c;那么就是下架超级快&#xff0c;有可能上午上…
阅读更多...
（一）万字详解G1垃圾收集器 —G1的设计目标是什么？G1的分区是什么？卡表的作用和工作原理？如何解决漏标问题？

（一）万字详解G1垃圾收集器 —G1的设计目标是什么？G1的分区是什么？卡表的作用和工作原理？如何解决漏标问题？

一、G1垃圾收集器简介 G1 GC&#xff08;Garbage-First Garbage Collector&#xff09;是一款先进的垃圾收集器&#xff0c;通过 -XX:UseG1GC 参数启用。它首次亮相于JDK 6u14版本&#xff0c;并在JDK 7u4中正式发布。对于熟悉JVM的开发者而言&#xff0c;G1已是一个广为人知的…
阅读更多...
html 获取浏览器地址栏参数

html 获取浏览器地址栏参数

例如&#xff1a; http://127.0.0.2/?agent_id143 我要获取agent_id 就很麻烦得去做字符串分割 解决方案&#xff1a; // 假设当前页面URL是: http://example.com/?param1value1&param2value2// 创建一个URLSearchParams对象 const params new URLSearchParams(wi…
阅读更多...
在线毫米(mm)到像素(px)换算器

在线毫米(mm)到像素(px)换算器

具体请前往&#xff1a;在线mm转px工具--将实际长度毫米(Millimeters)单位换算为像素(Pixels)单位
阅读更多...
打造高效舒适的气膜网球馆—轻空间

打造高效舒适的气膜网球馆—轻空间

气膜网球馆&#xff0c;作为现代运动设施的创新选择&#xff0c;其成本构成涵盖多个重要方面&#xff0c;确保为运动者提供最佳体验。 一、膜材选择 膜材是气膜网球馆的核心&#xff0c;品质不同直接影响成本。高品质膜材不仅增强了耐用性&#xff0c;也能有效阻挡外界气候影响…
阅读更多...
最新文章

Copyright @ 2022~2023