GPU Puzzles讲解（二）

 GPU-Puzzles项目是一个很棒的学习cuda编程的项目，可以让你学习到GPU编程和cuda核心并行编程的概念，通过一个个小问题让你理解cuda的编程和调用，创建共享显存空间，实现卷积和矩阵乘法等

https://github.com/srush/GPU-Puzzlesicon-default.png?t=O83Ahttps://github.com/srush/GPU-Puzzles

本文是接续我上一篇文章的讲解，深入分析几个比较困难的puzzles，讲解实现和优化原理

，GPU Puzzles讲解（一）-CSDN博客GPU Puzzles是一个很棒的cuda编程学习仓库，本文是对其中题目介绍和思量讲解https://blog.csdn.net/lijj0304/article/details/142737859GPU Puzzles是一个很棒的cuda编程学习仓库，本文是对其中题目介绍和思量讲解https://blog.csdn.net/lijj0304/article/details/142737859

 Puzzle 11 - 1D Convolution

Implement a kernel that computes a 1D convolution between a and b and stores it in out. You need to handle the general case. You only need 2 global reads and 1 global write per thread.

实现一维卷积的教学，这里要求控制全局读写次数。一开始很自然的会想到把待卷积的向量和卷积核内容分别存下来。同时考虑到共享内存空间一般和每块的线程数直接关联，而卷积需要乘周围的几个元素，我们需要额外存储边界的几个元素。不仅要在共享内存中存储卷积核的内容，又为了控制到2次的全局读写，这里用了一个trick：即利用不需要存卷积核时的索引去存边界的额外元素

def conv_spec(a, b):out = np.zeros(*a.shape)len = b.shape[0]for i in range(a.shape[0]):out[i] = sum([a[i + j] * b[j] for j in range(len) if i + j < a.shape[0]])return outMAX_CONV = 4
TPB = 8
TPB_MAX_CONV = TPB + MAX_CONV
def conv_test(cuda):def call(out, a, b, a_size, b_size) -> None:i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.xlocal_i = cuda.threadIdx.x# FILL ME IN (roughly 17 lines)shared_a = cuda.shared.array(TPB_MAX_CONV, numba.float32)shared_b = cuda.shared.array(MAX_CONV, numba.float32)if i < a_size:shared_a[local_i] = a[i]if local_i < b_size:shared_b[local_i] = b[local_i]else:local_j = local_i - b_sizeif i-b_size+TPB < a_size and local_j < b_size:shared_a[local_j+TPB] = a[i-b_size+TPB]cuda.syncthreads()total = 0.0for j in range(b_size):if i+j < a_size:total += shared_a[local_i+j] * shared_b[j]if i < a_size:out[i] = totalreturn call# Test 1SIZE = 6
CONV = 3
out = np.zeros(SIZE)
a = np.arange(SIZE)
b = np.arange(CONV)
problem = CudaProblem("1D Conv (Simple)",conv_test,[a, b],out,[SIZE, CONV],Coord(1, 1),Coord(TPB, 1),spec=conv_spec,
)
problem.show()

可视化效果

# 1D Conv (Simple)Score (Max Per Thread):|  Global Reads | Global Writes |  Shared Reads | Shared Writes ||             2 |             1 |             6 |             2 |

​

一维卷积Test2：向量长度超过TPB

out = np.zeros(15)
a = np.arange(15)
b = np.arange(4)
problem = CudaProblem("1D Conv (Full)",conv_test,[a, b],out,[15, 4],Coord(2, 1),Coord(TPB, 1),spec=conv_spec,
)
problem.show()

 可视化效果

# 1D Conv (Full)Score (Max Per Thread):|  Global Reads | Global Writes |  Shared Reads | Shared Writes ||             2 |             1 |             8 |             2 |

​

Puzzle 12 - Prefix Sum

Implement a kernel that computes a sum over a and stores it in out. If the size of a is greater than the block size, only store the sum of each block.

​

这个是利用共享内存来计算前缀和，减少访存次数。可以利用两两归并相加的原理，最终结果存储在共享内存块的第1个。利用TPB分割向量长度，输出的out是每一段TPB长度的向量和

TPB = 8
def sum_spec(a):out = np.zeros((a.shape[0] + TPB - 1) // TPB)for j, i in enumerate(range(0, a.shape[-1], TPB)):out[j] = a[i : i + TPB].sum()return outdef sum_test(cuda):def call(out, a, size: int) -> None:cache = cuda.shared.array(TPB, numba.float32)i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.xlocal_i = cuda.threadIdx.x# FILL ME IN (roughly 12 lines)if i < size:cache[local_i] = a[i]else:cache[local_i] = 0cuda.syncthreads()p = 2while (p <= TPB):if local_i % p == 0:cache[local_i] += cache[local_i+p/2]cuda.syncthreads()p *= 2if local_i == 0:out[cuda.blockIdx.x] = cache[local_i]return call# Test 1SIZE = 8
out = np.zeros(1)
inp = np.arange(SIZE)
problem = CudaProblem("Sum (Simple)",sum_test,[inp],out,[SIZE],Coord(1, 1),Coord(TPB, 1),spec=sum_spec,
)
problem.show()

可视化效果

# Sum (Simple)Score (Max Per Thread):|  Global Reads | Global Writes |  Shared Reads | Shared Writes ||             1 |             1 |             7 |             4 |

​

前缀和Test2：

SIZE = 15
out = np.zeros(2)
inp = np.arange(SIZE)
problem = CudaProblem("Sum (Full)",sum_test,[inp],out,[SIZE],Coord(2, 1),Coord(TPB, 1),spec=sum_spec,
)
problem.show()

可视化效果

# Sum (Full)Score (Max Per Thread):|  Global Reads | Global Writes |  Shared Reads | Shared Writes ||             1 |             1 |             7 |             4 |

 ​​

 Puzzle 13 - Axis Sum

Implement a kernel that computes a sum over each column of a and stores it in out.

其实就是前缀和的运用，这里加入了batch作为列标计，每个batch计算和，out尺寸和bath大小一致

TPB = 8
def sum_spec(a):out = np.zeros((a.shape[0], (a.shape[1] + TPB - 1) // TPB))for j, i in enumerate(range(0, a.shape[-1], TPB)):out[..., j] = a[..., i : i + TPB].sum(-1)return outdef axis_sum_test(cuda):def call(out, a, size: int) -> None:cache = cuda.shared.array(TPB, numba.float32)i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.xlocal_i = cuda.threadIdx.xbatch = cuda.blockIdx.y# FILL ME IN (roughly 12 lines)if i < size:cache[local_i] = a[batch, i]else:cache[local_i] = 0cuda.syncthreads() p = 2while (p <= TPB):if local_i % p == 0:cache[local_i] += cache[local_i+p/2]cuda.syncthreads()p *= 2if local_i == 0:out[batch, 0] += cache[local_i]return callBATCH = 4
SIZE = 6
out = np.zeros((BATCH, 1))
inp = np.arange(BATCH * SIZE).reshape((BATCH, SIZE))
problem = CudaProblem("Axis Sum",axis_sum_test,[inp],out,[SIZE],Coord(1, BATCH),Coord(TPB, 1),spec=sum_spec,
)
problem.show()

可视化效果

# Axis SumScore (Max Per Thread):|  Global Reads | Global Writes |  Shared Reads | Shared Writes ||             2 |             1 |             7 |             4 |

​

Puzzle 14 - Matrix Multiply!

Implement a kernel that multiplies square matrices a and b and stores the result in out.

矩阵乘法实现教学。第一直觉肯定就是利用二维的线程矩阵，循环行列遍历乘积求和。但是会有一个致命的问题就是线程矩阵比实际的矩阵小，而且想要控制全局读取的次数。这时就要利用好二维的共享空间，可以从对应行列出发，按行横向，列竖向顺序取部分的（TPB，TPB）矩阵块然后求乘积。从Test2的可视化图可以很直观看出这个思想

def matmul_spec(a, b):return a @ bTPB = 3
def mm_oneblock_test(cuda):def call(out, a, b, size: int) -> None:a_shared = cuda.shared.array((TPB, TPB), numba.float32)b_shared = cuda.shared.array((TPB, TPB), numba.float32)i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.xj = cuda.blockIdx.y * cuda.blockDim.y + cuda.threadIdx.ylocal_i = cuda.threadIdx.xlocal_j = cuda.threadIdx.y# FILL ME IN (roughly 14 lines)total = 0.0for k in range(0, size, TPB): if local_j + k < size and i < size:a_shared[local_i, local_j] = a[i, local_j+k]if local_i + k < size and j < size:b_shared[local_i, local_j] = b[local_i+k, j]cuda.syncthreads()for local_k in range(min(TPB, size-k)):total += a_shared[local_i, local_k] * b_shared[local_k, local_j]if i < size and j < size:out[i, j] = totalreturn call# Test 1SIZE = 2
out = np.zeros((SIZE, SIZE))
inp1 = np.arange(SIZE * SIZE).reshape((SIZE, SIZE))
inp2 = np.arange(SIZE * SIZE).reshape((SIZE, SIZE)).Tproblem = CudaProblem("Matmul (Simple)",mm_oneblock_test,[inp1, inp2],out,[SIZE],Coord(1, 1),Coord(TPB, TPB),spec=matmul_spec,
)
problem.show(sparse=True)

可视化效果

# Matmul (Simple)Score (Max Per Thread):|  Global Reads | Global Writes |  Shared Reads | Shared Writes ||             2 |             1 |             4 |             2 |

​

Test2大于线程矩阵的矩阵乘法：

SIZE = 8
out = np.zeros((SIZE, SIZE))
inp1 = np.arange(SIZE * SIZE).reshape((SIZE, SIZE))
inp2 = np.arange(SIZE * SIZE).reshape((SIZE, SIZE)).Tproblem = CudaProblem("Matmul (Full)",mm_oneblock_test,[inp1, inp2],out,[SIZE],Coord(3, 3),Coord(TPB, TPB),spec=matmul_spec,
)
problem.show(sparse=True)

可视化效果

# Matmul (Full)Score (Max Per Thread):|  Global Reads | Global Writes |  Shared Reads | Shared Writes ||             6 |             1 |            16 |             6 |

​

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

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

相关文章

【正则表达式】粗浅学习

【正则表达式】粗浅学习

文章目录 1 背景介绍1.1 为什么要学习正则表达式&#xff1f;1.2 什么是正则表达式&#xff1f;1.3 正则表达式可以做什么&#xff1f; 2 基本语法3 典型案例3.1 正整数3.2 IPv4类型的IP地址3.2.1 点分10进制表示格式 1 背景介绍 1.1 为什么要学习正则表达式&#xff1f; 日常…
阅读更多...
面试系列-分组Tire树匹配算法

面试系列-分组Tire树匹配算法

自己写的分组Tire树匹配算法&#xff0c;该算法用于云南省人工智能重点实验室与云南电网合作项目&#xff08;云南电网敏感信息识别系统&#xff09;&#xff0c;用于快速匹配文本将项目中数据算法抽离出来&#xff0c;特此分享&#xff01;&#xff01;!可以实现动态的插入、删…
阅读更多...
NVIDIA Hopper 架构深入

NVIDIA Hopper 架构深入

在 2022 年 NVIDIA GTC 主题演讲中，NVIDIA 首席执行官黄仁勋介绍了基于全新 NVIDIA Hopper GPU 架构的全新 NVIDIA H100 Tensor Core GPU。 文章目录 前言一、NVIDIA H100 Tensor Core GPU 简介二、NVIDIA H100 GPU 主要功能概述1. 新的流式多处理器 （SM） 具有许多性能和效率…
阅读更多...
leetcode58:最后一个单词的长度

leetcode58:最后一个单词的长度

给你一个字符串 s&#xff0c;由若干单词组成&#xff0c;单词前后用一些空格字符隔开。返回字符串中 最后一个 单词的长度。 单词 是指仅由字母组成、不包含任何空格字符的最大 子字符串 。 示例 1&#xff1a; 输入&#xff1a;s "Hello World" 输出&#xff…
阅读更多...
浅谈汽车智能座舱如何实现多通道音频

浅谈汽车智能座舱如何实现多通道音频

一、引言 随着汽车智能座舱的功能迭代发展&#xff0c;传统的 4 通道、6 通道、8 通道等音响系统难以在满足驾驶场景的需求&#xff0c;未来对于智能座舱音频质量和通道数会越来越高。接下来本文将浅析目前智能座舱如何实现音频功放&#xff0c;以及如何实现多路音频功放方案。…
阅读更多...
iOS 多次获取图片主题色不一样

iOS 多次获取图片主题色不一样

一个需求中&#xff0c;要求获取图片的主题色 代码如下 -(void)kk_getImage:(UIImage *)image fetchthemeColor:(void(^)(UIColor *color))callBack {dispatch_async(dispatch_get_global_queue(DISPATCH_QUEUE_PRIORITY_DEFAULT, 0), ^{// 第一步 先把图片缩小 加快计算速度.…
阅读更多...
C语言文件操作(上)(27)

C语言文件操作(上)(27)

文章目录 前言一、为什么要用文件&#xff1f;二、什么是文件&#xff1f;程序文件数据文件文件名文件类型文件缓冲区文件指针 三、流流的概念标准流 总结 前言 C语言可以直接操作文件&#xff0c;如果你是第一次听说这个特性&#xff0c;可能会眼前一亮&#xff0c;感到惊奇  …
阅读更多...
MongoDB的安装与增删改查基本操作

MongoDB的安装与增删改查基本操作

MongoDB是一种非关系型数据库，是NoSQL语言，但是又是最接近关系型数据库的。内部存储不是表结构，但是可以对数据进行表结构的操作。 一、安装 在官网：Download MongoDB Community Server | MongoDB下载系统对应的版本进行安装即可 二、编辑器 在安装MongoDB后会自带一个编…
阅读更多...
图片格式入门

图片格式入门

主要参考资料&#xff1a; 常见的图片格式介绍: https://blog.csdn.net/cnds123/article/details/127165291 目录 像素图与矢量图像素图&#xff08;pixel image&#xff09;矢量图&#xff08;Vector graphics&#xff09; 像素图与矢量图 像素图&#xff08;pixel image&…
阅读更多...
D3.js中国地图可视化

D3.js中国地图可视化

1、项目介绍 该项目来自Github&#xff0c;基于D3.js中国地图可视化。 D3.js is a JavaScript library for manipulating documents based on data. It uses HTML, SVG, and CSS to display data. The full name of D3 is "Data-Driven Documents," which means it a…
阅读更多...
Spring Boot项目使用MyBatis Plus的详细步骤

Spring Boot项目使用MyBatis Plus的详细步骤

在Spring Boot项目中使用MyBatis Plus&#xff0c;可以极大地简化数据库操作&#xff0c;提高开发效率。以下是在Spring Boot项目中集成和使用MyBatis Plus的详细步骤&#xff1a; 一、环境准备 确保已安装Java和Spring Boot&#xff1a;MyBatis Plus是基于Java和Spring Boot…
阅读更多...
话术挂断之后是否处理事件

话术挂断之后是否处理事件

文章目录 前言联系我们解决方案方案一方案二 前言 流程&#xff1a;自动外呼进入机器人话术。问题&#xff1a;在机器人放音时用户挂断后&#xff0c;话术还会继续匹配流程&#xff0c;如果匹配上的是放音节点&#xff0c;还会进行放音&#xff0c;那么在数据库表conversation…
阅读更多...
Redis 缓存策略详解：提升性能的四种常见模式

Redis 缓存策略详解：提升性能的四种常见模式

在现代分布式系统中&#xff0c;缓存是提升性能和减轻数据库负载的关键组件。Redis 作为一种高性能的内存数据库&#xff0c;被广泛应用于缓存层。本文将深入探讨几种常用的 Redis 缓存策略&#xff0c;包括旁路缓存模式&#xff08;Cache-Aside Pattern&#xff09;、读穿透模…
阅读更多...
windows中C++调用dll文件的两种方式，QT QLibrary和Windows API

windows中C++调用dll文件的两种方式，QT QLibrary和Windows API

假设DLL文件名为 test.dll&#xff0c;函数 int add(int a.int b) 1.使用Windows API 的方式 #include <windows.h>extern "C" __declspec(dllexport) int add(int a, int b); // 定义函数typedef int (*AddFunction)(int, int); // 定义函数指针类型HMODULE …
阅读更多...
阿里云 SAE Web：百毫秒高弹性的实时事件中心的架构和挑战

阿里云 SAE Web：百毫秒高弹性的实时事件中心的架构和挑战

作者&#xff1a;胡志广(独鳌) 背景 Serverless 应用引擎 SAE 事件中心主要面向早期的 SAE 控制台只有针对于应用维度的事件&#xff0c;这个事件是 K8s 原生的事件&#xff0c;其实绝大多数的用户并不会关心&#xff0c;同时也可能看不懂。而事件中心&#xff0c;是希望能够…
阅读更多...
SpringBoot MyBatis连接数据库设置了encoding=utf-8还是不能用中文来查询

SpringBoot MyBatis连接数据库设置了encoding=utf-8还是不能用中文来查询

properties的MySQL连接时已经指定了字符编码格式&#xff1a; url: jdbc:mysql://localhost:3306/sky_take_out?useUnicodetrue&characterEncodingutf-8使用MyBatis查询&#xff0c;带有中文参数&#xff0c;查询出的内容为空。 执行的语句为&#xff1a; <select id&…
阅读更多...
一个月冲刺软考——病毒与木马的了解、认证与加密、加密技术的分类

一个月冲刺软考——病毒与木马的了解、认证与加密、加密技术的分类

目录 1.计算机病毒与木马 2.加密方式的分类及特点 3.选择路由的策略 4.DoS攻击(拒绝服务) 5.认证与加密 6.报文摘要算法 7.网络攻击的类别 8.安全的分类及举例 9.加密技术的分类 10.什么是PKI体制&#xff1f;有哪些步骤&#xff1f; 1.计算机病毒与木马 1.1病毒&am…
阅读更多...
Tensorflow2.0

Tensorflow2.0

Tensorflow2.0 有深度学习基础的建议直接看class3 class1 介绍 人工智能3学派 行为主义:基于控制论&#xff0c;构建感知-动作控制系统。(控制论&#xff0c;如平衡、行走、避障等自适应控制系统) 符号主义:基于算数逻辑表达式&#xff0c;求解问题时先把问题描述为表达式…
阅读更多...
开源跨平台三维模型轻量化软件osgGISPlugins-1、简介

开源跨平台三维模型轻量化软件osgGISPlugins-1、简介

下一篇文章&#xff1a;开源跨平台三维模型轻量化软件osgGISPlugins-2、如何编译 1、仓库地址 github地址&#xff1a;https://github.com/newpeople123/osgGISPlugins gitee地址&#xff1a;https://gitee.com/wtyhz/osg-gis-plugins 2、项目简介 osg引擎的gis插件&#x…
阅读更多...
【Kubernetes】常见面试题汇总（五十三）

【Kubernetes】常见面试题汇总（五十三）

目录 118. pod 状态为 ErrlmagePull &#xff1f; 119.探测存活 pod 状态为 CrashLoopBackOff &#xff1f; 特别说明&#xff1a; 题目 1-68 属于【Kubernetes】的常规概念题&#xff0c;即 “ 汇总&#xff08;一&#xff09;~&#xff08;二十二&#xff09;” 。…
阅读更多...
最新文章

Copyright @ 2022~2023