PPMP_char3
PMPP char3 – Multidimensional grids and data
五一过后,有些工作要赶,抽出时间更新一下。这一章基本都熟练掌握,在做习题过程中有一些思考。这里涉及到了一点点GEMM(矩阵乘),GEMM有太多可深挖的了,推荐一篇博客How to Optimize a CUDA Matmul Kernel for cuBLAS-like Performance: a Worklog (siboehm.com)。另外,我还发现上一篇博客,有写错的地方,这篇博客的末尾做了一下勘误。这里记录我的个人理解,有不正确的地方,欢迎留言或者私信讨论。
课后习题
- In this chapter we implemented a matrix multiplication kernel that has each
thread produce one output matrix element. In this question, you will
implement different matrix-matrix multiplication kernels and compare them.
a. Write a kernel that has each thread produce one output matrix row. Fill in
the execution configuration parameters for the design.
b. Write a kernel that has each thread produce one output matrix column. Fill
in the execution configuration parameters for the design.
c. Analyze the pros and cons of each of the two kernel designs.
答案:
a 部分 (来自大模型)
__global__ void matrixMulRow(float *A, float *B, float *C, int m, int n, int k) {int row = blockIdx.y * blockDim.y + threadIdx.y;if (row < m) {for (int col = 0; col < k; ++col) {float sum = 0;for (int i = 0; i < n; ++i) {sum += A[row * n + i] * B[i * k + col];}C[row * k + col] = sum;}}
}
dim3 threadsPerBlock(1, 256);
dim3 blocksPerGrid(1, (m + threadsPerBlock.y - 1) / threadsPerBlock.y);
matrixMulRow<<<blocksPerGrid, threadsPerBlock>>>(A, B, C, m, n, k);
b部分 (来自大模型)
__global__ void matrixMulCol(float *A, float *B, float *C, int m, int n, int k) {int col = blockIdx.x * blockDim.x + threadIdx.x;if (col < k) {for (int row = 0; row < m; ++row) {float sum = 0;for (int i = 0; i < n; ++i) {sum += A[row * n + i] * B[i * k + col];}C[row * k + col] = sum;}}
}
dim3 threadsPerBlock(256, 1);
dim3 blocksPerGrid((k + threadsPerBlock.x - 1) / threadsPerBlock.x, 1);
matrixMulCol<<<blocksPerGrid, threadsPerBlock>>>(A, B, C, m, n, k);
c部分
假设A、B、C都是行主序;这里也不使用共享内存。一次读入一个缓存行,访问一行数据的话,存在访存局部性,需要用到的数据就在缓存中;访问一列数据的话,不存在访存局部性,需要用到的数据就不在缓存中,需要再读一个缓存行。假设一个缓存行时64B,对应16个float32。
访问A | 访问B | 访问C | |
---|---|---|---|
一个线程处理C中的一行 | 连续访问,只访问一行。访存次数=K/16 | 不连续访问,一次访问一列,访存次数=K;总共要访问N列,访问次数=N*K | 连续访问,只访问一行。访存次数=K/16 |
一个线程处理C中的一列 | 不连续访问,只访问一列。访存次数=M | 不连续访问,一次访问一列,访存次数=K;总共要访问N列,访问次数=N*K | 不连续访问,只访问一列。访存次数=M |
如果访存数据量较小,也就是a那种,可以直接放到寄存器中,这样访存cycle更短。
如果访存数据量较大,reg的容量就不够用,需要一部分存到L1上面,甚至L1也装不下,要从global里读取。涉及多个存储数据一致性、data hazard的问题,效率就很低。
- A matrix-vector multiplication takes an input matrix B and a vector C and
produces one output vector A. Each element of the output vector A is the dot
product of one row of the input matrix B and C, that is, A[i] = ΣB[i][j] * C[j].
For simplicity we will handle only square matrices whose elements are single-
precision floating-point numbers. Write a matrix-vector multiplication kernel and
the host stub function that can be called with four parameters: pointer to the output
matrix, pointer to the input matrix, pointer to the input vector, and the number of
elements in each dimension. Use one thread to calculate an output vector element.
答案:
其实,这道题就是参考gemm,实现gemv。就是把输出C的一个变量(C[i])映射到一个线程上,这个线程遍历j个变量,再各自点乘+求和(也就是j维度上做规约)。
以下是大模型写出来的程序,我看了下没啥问题,测试了下也是ok的。输入为A和B,输出为C。
#include <iostream>
#include <vector>
#include <random>
#include <cuda_runtime.h>#define CHECK(call) \
{ \const cudaError_t error = call; \if (error != cudaSuccess) { \std::cout << "Error: " << __FILE__ << ":" << __LINE__ << ", " << cudaGetErrorString(error) << std::endl; \exit(1); \} \
}__global__ void matrixVectorMultiply(const float* A, const float* B, float* C, int rows, int cols) {int i = blockIdx.x * blockDim.x + threadIdx.x;if (i < rows) {float sum = 0.0f;for (int j = 0; j < cols; ++j) {sum += A[j] * B[i * cols + j];}C[i] = sum;}
}int main() {int rows = 1024;int cols = 768;std::vector<float> hostA(cols);std::vector<float> hostB(rows * cols);std::vector<float> hostC(rows);std::vector<float> resultC(rows);// 初始化输入数据std::random_device rd;std::mt19937 gen(rd());std::uniform_real_distribution<float> dis(-1.0, 1.0);for (int j = 0; j < cols; ++j) {hostA[j] = dis(gen);}for (int i = 0; i < rows; ++i) {for (int j = 0; j < cols; ++j) {hostB[i * cols + j] = dis(gen);}}// 分配设备内存float* deviceA;float* deviceB;float* deviceC;CHECK(cudaMalloc(&deviceA, cols * sizeof(float)));CHECK(cudaMalloc(&deviceB, rows * cols * sizeof(float)));CHECK(cudaMalloc(&deviceC, rows * sizeof(float)));// 将输入数据从主机复制到设备CHECK(cudaMemcpy(deviceA, hostA.data(), cols * sizeof(float), cudaMemcpyHostToDevice));CHECK(cudaMemcpy(deviceB, hostB.data(), rows * cols * sizeof(float), cudaMemcpyHostToDevice));// 启动内核int threadsPerBlock = 256;int blocksPerGrid = (rows + threadsPerBlock - 1) / threadsPerBlock;matrixVectorMultiply<<<blocksPerGrid, threadsPerBlock>>>(deviceA, deviceB, deviceC, rows, cols);CHECK(cudaGetLastError());// 将输出数据从设备复制到主机CHECK(cudaMemcpy(hostC.data(), deviceC, rows * sizeof(float), cudaMemcpyDeviceToHost));// 验证结果正确性for (int i = 0; i < rows; ++i) {float sum = 0.0f;for (int j = 0; j < cols; ++j) {sum += hostA[j] * hostB[i * cols + j];}resultC[i] = sum;}for (int i = 0; i < rows; ++i) {if (std::abs(hostC[i] - resultC[i]) > 1e-5) {std::cout << "Result verification failed at index " << i << std::endl;return 1;}}std::cout << "Result verification passed" << std::endl;// 释放设备内存CHECK(cudaFree(deviceA));CHECK(cudaFree(deviceB));CHECK(cudaFree(deviceC));return 0;
}
- Consider the following CUDA kernel and the corresponding host function that
calls it:a. What is the number of threads per block?
b. What is the number of threads in the grid?c. What is the number of blocks in the grid?
d. What is the number of threads that execute the code on line 05?
答案:a、16*32=512;b、48640;c、95;d、45000
- Consider a 2D matrix with a width of 400 and a height of 500. The matrix is
stored as a one-dimensional array. Specify the array index of the matrix
element at row 20 and column 10:
a. If the matrix is stored in row-major order.
b. If the matrix is stored in column-major order.
答案:a、20*400+10=8010
b、10*500+20=5020
Consider a 3D tensor with a width of 400, a height of 500, and a depth of 300
The tensor is stored as a one-dimensional array in row-major order.
Specify the array index of the tensor element at x 5 10, y 5 20, and z 5 5.
答案: 5 * (400 * 500) + 10 * 400 + 5 = 1004005
勘误
上一篇博客中,这个图里面的d_a的类型写错了,应该是int* d_a,而不是void* d_a。传入函数的是(void**)&d_a,函数结束后,d_a还是int*类型。
相关文章:

PPMP_char3
PMPP char3 – Multidimensional grids and data 五一过后,有些工作要赶,抽出时间更新一下。这一章基本都熟练掌握,在做习题过程中有一些思考。这里涉及到了一点点GEMM(矩阵乘),GEMM有太多可深挖的了&a…...
VulkanSDK Demos vkcube 编译失败
操作系统: Windows 11 23H2 Vulkan 版本: 1.3.2.280.0 Visual Studio 版本: 2022 在VulkanSDK/Demos目录下存在一个demo solution,其中包含两个project, vkcube和vkcubepp,两个分别为C语言和C写的示例程序, 但是直接编译这两个project时会编译失败,报了以下错误: fatal err…...

(二)Jetpack Compose 布局模型
前文回顾 (一)Jetpack Compose 从入门到会写-CSDN博客 首先让我们回顾一下上一篇文章中里提到过几个问题: ComposeView的层级关系,互相嵌套存在的问题? 为什么Compose可以实现只测量一次? ComposeView和…...
【Oracle impdp导入dmp文件(windows)】
Oracle impdp导入dmp文件(windows) 1、连接数据库2、创建与导出的模式相同名称的用户WIRELESS2,并赋予权限3、创建directory 的物理目录f:\radio\dmp,并把.dmp文件放进去4、连接新用户WIRELESS25、创建表空间的物理目录F:\radio\t…...

代数结构:5、格与布尔代数
16.1 偏序与格 偏序集:设P是集合,P上的二元关系“≤”满足以下三个条件,则称“≤”是P上的偏序关系(或部分序关系) (1)自反性:a≤a,∀a∈P; (2…...

如何使用DEEPL免费翻译PDF
如何使用DEEPL免费翻译PDF 安装DEEPL取消PDF限制 安装DEEPL 安装教程比较多,这里不重复。 把英文pdf拖进去,点翻译,在下面的框中有已经翻译完毕的文档。 但是存在两个问题 问题1:这些文档是加密的。 问题2:带有DeepL标…...
Spring-全面详解
Spring,就像是软件开发界的一个超级英雄,它让编写Java程序变得更简单、更灵活。想象一下,如果你要盖一栋大楼,Spring就是那个提供各种工具、框架和最佳实践的建筑大师,帮助你高效、优雅地搭建起整个项目。 Spring是啥&…...

QT自适应界面 处理高DPI 缩放比界面乱问题
1.pro文件添加 必须添加要不找不到 QT版本需要 5。4 以上才支持 QT widgets 2.main界面提前处理 // 1. 全局缩放使能QApplication::setAttribute(Qt::AA_EnableHighDpiScaling, true);// 2. 适配非整数倍缩放QGuiApplication::setHighDpiScaleFactorRoundingPolicy(Qt::High…...

序列到序列模型在语言识别Speech Applications中的应用 Transformer应用于TTS Transformer应用于ASR 端到端RNN
序列到序列模型在语言识别Speech Applications中的应用 A Comparative Study on Transformer vs RNN in Speech Applications 序列到序列(Seq2Seq)模型在语音识别(Speech Applications)中有重要的应用。虽然Seq2Seq模型最初是为了解决自然语言处理中的序列生成问题而设计的…...

【Linux】- Linux环境变量[8]
目录 环境变量 $符号 自行设置环境变量 环境变量 环境变量是操作系统(Windows、Linux、Mac)在运行的时候,记录的一些关键性信息,用以辅助系统运行。在Linux系统中执行:env命令即可查看当前系统中记录的环境变量。 …...

前端笔记-day04
文章目录 01-后代选择器02-子代选择器03-并集选择器04-交集选择器05-伪类选择器06-拓展-超链接伪类07-CSS特性-继承性08-CSS特性-层叠性09-CSS特性-优先级11-Emmet写法12-背景图13-背景图平铺方式14-背景图位置15-背景图缩放16-背景图固定17-background属性18-显示模式19-显示模…...

计算机字符集产生的历史与乱码
你好,我是 shengjk1,多年大厂经验,努力构建 通俗易懂的、好玩的编程语言教程。 欢迎关注!你会有如下收益: 了解大厂经验拥有和大厂相匹配的技术等 希望看什么,评论或者私信告诉我! 文章目录 一…...

Rerank进一步提升RAG效果
RAG & Rerank 目前大模型应用中,RAG(Retrieval Augmented Generation,检索增强生成)是一种在对话(QA)场景下最主要的应用形式,它主要解决大模型的知识存储和更新问题。 简述RAG without R…...

使用train.py----yolov7
准备工作 在训练之前,数据集的工作和配置环境的工作要做好 数据集:看这里划分数据集,训练自己的数据集。_划分数据集后如何训练-CSDN博客 划分数据集2,详细说明-CSDN博客 配置环境看这里 从0开始配置环境-yolov7_gpu0是inter g…...

机器学习第37周周报 GGNN
文章目录 week37 GGNN摘要Abstract一、文献阅读1. 题目2. abstract3. 网络架构3.1 数据处理部分3.2 门控图神经网络3.3 掩码操作 4. 文献解读4.1 Introduction4.2 创新点4.3 实验过程4.3.1 传感器设置策略4.3.2 数据集4.3.3 实验设置4.3.4 模型参数设置4.3.5 实验结果 5. 结论 …...

Baidu Comate:释放编码潜能,革新软件开发
Baidu Comate Baidu Comate,智能代码助手,凭借着文心大模型的强大支撑,结合了百度多年的编程实战数据和丰富的开源资源,形成了一款崭新的编码辅助利器。它不仅具备着高智能、多场景、价值创造的特质,更可广泛应用于各…...

MATLAB的Bar3函数调节渐变色(内附渐变色库.mat及.m文件免费下载链接)
一. colormap函数 可以使用colormap函数: t1[281.1,584.6, 884.3,1182.9,1485.2; 291.6,592.6,896,1197.75,1497.33; 293.8,596.4,898.6,1204.4,1506.4; 295.8,598,904.4,1209.0,1514.6];bar3(t1,1) set(gca,XTickLabel,{300,600,900,1200,1500},FontSize,10) set…...
使用 TensorFlow.js 和 OffscreenCanvas 实现实时防挡脸弹幕
首先,要理解我们的目标,我们将实时获取视频中的面部区域并将其周围的内容转为不透明以制造出弹幕的“遮挡效应”。 步骤一:环境准备 我们将使用 TensorFlow.js 的 Body-segmentation 库来完成面部识别部分,并使用 OffscreenCanv…...

【计算机网络篇】数据链路层(10)在物理层扩展以太网
文章目录 🍔扩展站点与集线器之间的距离🛸扩展共享式以太网的覆盖范围和站点数量 🍔扩展站点与集线器之间的距离 🛸扩展共享式以太网的覆盖范围和站点数量 以太网集线器一般具有8~32个接口,如果要连接的站点数量超过了…...

conan2 基础入门(03)-使用(msvc为例)
conan2 基础入门(03)-使用(msvc为例) 文章目录 conan2 基础入门(03)-使用(msvc为例)⭐准备生成profile文件预备文件和Code ⭐使用指令预览正确执行结果可能出现的问题 ⭐具体讲解conanconanfile.txt执行 install cmakeCMakeLists.txt生成项目构建 END ⭐准备 在阅读和学习本文…...
React 第五十五节 Router 中 useAsyncError的使用详解
前言 useAsyncError 是 React Router v6.4 引入的一个钩子,用于处理异步操作(如数据加载)中的错误。下面我将详细解释其用途并提供代码示例。 一、useAsyncError 用途 处理异步错误:捕获在 loader 或 action 中发生的异步错误替…...
多场景 OkHttpClient 管理器 - Android 网络通信解决方案
下面是一个完整的 Android 实现,展示如何创建和管理多个 OkHttpClient 实例,分别用于长连接、普通 HTTP 请求和文件下载场景。 <?xml version"1.0" encoding"utf-8"?> <LinearLayout xmlns:android"http://schemas…...

P3 QT项目----记事本(3.8)
3.8 记事本项目总结 项目源码 1.main.cpp #include "widget.h" #include <QApplication> int main(int argc, char *argv[]) {QApplication a(argc, argv);Widget w;w.show();return a.exec(); } 2.widget.cpp #include "widget.h" #include &q…...

DIY|Mac 搭建 ESP-IDF 开发环境及编译小智 AI
前一阵子在百度 AI 开发者大会上,看到基于小智 AI DIY 玩具的演示,感觉有点意思,想着自己也来试试。 如果只是想烧录现成的固件,乐鑫官方除了提供了 Windows 版本的 Flash 下载工具 之外,还提供了基于网页版的 ESP LA…...
3403. 从盒子中找出字典序最大的字符串 I
3403. 从盒子中找出字典序最大的字符串 I 题目链接:3403. 从盒子中找出字典序最大的字符串 I 代码如下: class Solution { public:string answerString(string word, int numFriends) {if (numFriends 1) {return word;}string res;for (int i 0;i &…...

全志A40i android7.1 调试信息打印串口由uart0改为uart3
一,概述 1. 目的 将调试信息打印串口由uart0改为uart3。 2. 版本信息 Uboot版本:2014.07; Kernel版本:Linux-3.10; 二,Uboot 1. sys_config.fex改动 使能uart3(TX:PH00 RX:PH01),并让boo…...

图表类系列各种样式PPT模版分享
图标图表系列PPT模版,柱状图PPT模版,线状图PPT模版,折线图PPT模版,饼状图PPT模版,雷达图PPT模版,树状图PPT模版 图表类系列各种样式PPT模版分享:图表系列PPT模板https://pan.quark.cn/s/20d40aa…...
代理篇12|深入理解 Vite中的Proxy接口代理配置
在前端开发中,常常会遇到 跨域请求接口 的情况。为了解决这个问题,Vite 和 Webpack 都提供了 proxy 代理功能,用于将本地开发请求转发到后端服务器。 什么是代理(proxy)? 代理是在开发过程中,前端项目通过开发服务器,将指定的请求“转发”到真实的后端服务器,从而绕…...

九天毕昇深度学习平台 | 如何安装库?
pip install 库名 -i https://pypi.tuna.tsinghua.edu.cn/simple --user 举个例子: 报错 ModuleNotFoundError: No module named torch 那么我需要安装 torch pip install torch -i https://pypi.tuna.tsinghua.edu.cn/simple --user pip install 库名&#x…...

HarmonyOS运动开发:如何用mpchart绘制运动配速图表
##鸿蒙核心技术##运动开发##Sensor Service Kit(传感器服务)# 前言 在运动类应用中,运动数据的可视化是提升用户体验的重要环节。通过直观的图表展示运动过程中的关键数据,如配速、距离、卡路里消耗等,用户可以更清晰…...