当前位置: 首页 > news >正文

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)。另外,我还发现上一篇博客,有写错的地方,这篇博客的末尾做了一下勘误。这里记录我的个人理解,有不正确的地方,欢迎留言或者私信讨论。

课后习题

  1. 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的问题,效率就很低。

  1. 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;
}
  1. 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

  1. 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

  1. 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 ​ 五一过后&#xff0c;有些工作要赶&#xff0c;抽出时间更新一下。这一章基本都熟练掌握&#xff0c;在做习题过程中有一些思考。这里涉及到了一点点GEMM&#xff08;矩阵乘&#xff09;&#xff0c;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 布局模型

前文回顾 &#xff08;一&#xff09;Jetpack Compose 从入门到会写-CSDN博客 首先让我们回顾一下上一篇文章中里提到过几个问题&#xff1a; ComposeView的层级关系&#xff0c;互相嵌套存在的问题&#xff1f; 为什么Compose可以实现只测量一次&#xff1f; ComposeView和…...

【Oracle impdp导入dmp文件(windows)】

Oracle impdp导入dmp文件&#xff08;windows&#xff09; 1、连接数据库2、创建与导出的模式相同名称的用户WIRELESS2&#xff0c;并赋予权限3、创建directory 的物理目录f:\radio\dmp&#xff0c;并把.dmp文件放进去4、连接新用户WIRELESS25、创建表空间的物理目录F:\radio\t…...

代数结构:5、格与布尔代数

16.1 偏序与格 偏序集&#xff1a;设P是集合&#xff0c;P上的二元关系“≤”满足以下三个条件&#xff0c;则称“≤”是P上的偏序关系&#xff08;或部分序关系&#xff09; &#xff08;1&#xff09;自反性&#xff1a;a≤a&#xff0c;∀a∈P&#xff1b; &#xff08;2…...

如何使用DEEPL免费翻译PDF

如何使用DEEPL免费翻译PDF 安装DEEPL取消PDF限制 安装DEEPL 安装教程比较多&#xff0c;这里不重复。 把英文pdf拖进去&#xff0c;点翻译&#xff0c;在下面的框中有已经翻译完毕的文档。 但是存在两个问题 问题1&#xff1a;这些文档是加密的。 问题2&#xff1a;带有DeepL标…...

Spring-全面详解

Spring&#xff0c;就像是软件开发界的一个超级英雄&#xff0c;它让编写Java程序变得更简单、更灵活。想象一下&#xff0c;如果你要盖一栋大楼&#xff0c;Spring就是那个提供各种工具、框架和最佳实践的建筑大师&#xff0c;帮助你高效、优雅地搭建起整个项目。 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]

目录 环境变量 $符号 自行设置环境变量 环境变量 环境变量是操作系统&#xff08;Windows、Linux、Mac&#xff09;在运行的时候&#xff0c;记录的一些关键性信息&#xff0c;用以辅助系统运行。在Linux系统中执行&#xff1a;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-显示模…...

计算机字符集产生的历史与乱码

你好&#xff0c;我是 shengjk1&#xff0c;多年大厂经验&#xff0c;努力构建 通俗易懂的、好玩的编程语言教程。 欢迎关注&#xff01;你会有如下收益&#xff1a; 了解大厂经验拥有和大厂相匹配的技术等 希望看什么&#xff0c;评论或者私信告诉我&#xff01; 文章目录 一…...

Rerank进一步提升RAG效果

RAG & Rerank 目前大模型应用中&#xff0c;RAG&#xff08;Retrieval Augmented Generation&#xff0c;检索增强生成&#xff09;是一种在对话&#xff08;QA&#xff09;场景下最主要的应用形式&#xff0c;它主要解决大模型的知识存储和更新问题。 简述RAG without R…...

使用train.py----yolov7

准备工作 在训练之前&#xff0c;数据集的工作和配置环境的工作要做好 数据集&#xff1a;看这里划分数据集&#xff0c;训练自己的数据集。_划分数据集后如何训练-CSDN博客 划分数据集2&#xff0c;详细说明-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&#xff0c;智能代码助手&#xff0c;凭借着文心大模型的强大支撑&#xff0c;结合了百度多年的编程实战数据和丰富的开源资源&#xff0c;形成了一款崭新的编码辅助利器。它不仅具备着高智能、多场景、价值创造的特质&#xff0c;更可广泛应用于各…...

MATLAB的Bar3函数调节渐变色(内附渐变色库.mat及.m文件免费下载链接)

一. colormap函数 可以使用colormap函数&#xff1a; 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 实现实时防挡脸弹幕

首先&#xff0c;要理解我们的目标&#xff0c;我们将实时获取视频中的面部区域并将其周围的内容转为不透明以制造出弹幕的“遮挡效应”。 步骤一&#xff1a;环境准备 我们将使用 TensorFlow.js 的 Body-segmentation 库来完成面部识别部分&#xff0c;并使用 OffscreenCanv…...

【计算机网络篇】数据链路层(10)在物理层扩展以太网

文章目录 &#x1f354;扩展站点与集线器之间的距离&#x1f6f8;扩展共享式以太网的覆盖范围和站点数量 &#x1f354;扩展站点与集线器之间的距离 &#x1f6f8;扩展共享式以太网的覆盖范围和站点数量 以太网集线器一般具有8~32个接口&#xff0c;如果要连接的站点数量超过了…...

conan2 基础入门(03)-使用(msvc为例)

conan2 基础入门(03)-使用(msvc为例) 文章目录 conan2 基础入门(03)-使用(msvc为例)⭐准备生成profile文件预备文件和Code ⭐使用指令预览正确执行结果可能出现的问题 ⭐具体讲解conanconanfile.txt执行 install cmakeCMakeLists.txt生成项目构建 END ⭐准备 在阅读和学习本文…...

Linux 文件类型,目录与路径,文件与目录管理

文件类型 后面的字符表示文件类型标志 普通文件&#xff1a;-&#xff08;纯文本文件&#xff0c;二进制文件&#xff0c;数据格式文件&#xff09; 如文本文件、图片、程序文件等。 目录文件&#xff1a;d&#xff08;directory&#xff09; 用来存放其他文件或子目录。 设备…...

React Native 导航系统实战(React Navigation)

导航系统实战&#xff08;React Navigation&#xff09; React Navigation 是 React Native 应用中最常用的导航库之一&#xff0c;它提供了多种导航模式&#xff0c;如堆栈导航&#xff08;Stack Navigator&#xff09;、标签导航&#xff08;Tab Navigator&#xff09;和抽屉…...

【机器视觉】单目测距——运动结构恢复

ps&#xff1a;图是随便找的&#xff0c;为了凑个封面 前言 在前面对光流法进行进一步改进&#xff0c;希望将2D光流推广至3D场景流时&#xff0c;发现2D转3D过程中存在尺度歧义问题&#xff0c;需要补全摄像头拍摄图像中缺失的深度信息&#xff0c;否则解空间不收敛&#xf…...

Redis数据倾斜问题解决

Redis 数据倾斜问题解析与解决方案 什么是 Redis 数据倾斜 Redis 数据倾斜指的是在 Redis 集群中&#xff0c;部分节点存储的数据量或访问量远高于其他节点&#xff0c;导致这些节点负载过高&#xff0c;影响整体性能。 数据倾斜的主要表现 部分节点内存使用率远高于其他节…...

如何理解 IP 数据报中的 TTL?

目录 前言理解 前言 面试灵魂一问&#xff1a;说说对 IP 数据报中 TTL 的理解&#xff1f;我们都知道&#xff0c;IP 数据报由首部和数据两部分组成&#xff0c;首部又分为两部分&#xff1a;固定部分和可变部分&#xff0c;共占 20 字节&#xff0c;而即将讨论的 TTL 就位于首…...

蓝桥杯 冶炼金属

原题目链接 &#x1f527; 冶炼金属转换率推测题解 &#x1f4dc; 原题描述 小蓝有一个神奇的炉子用于将普通金属 O O O 冶炼成为一种特殊金属 X X X。这个炉子有一个属性叫转换率 V V V&#xff0c;是一个正整数&#xff0c;表示每 V V V 个普通金属 O O O 可以冶炼出 …...

Java毕业设计:WML信息查询与后端信息发布系统开发

JAVAWML信息查询与后端信息发布系统实现 一、系统概述 本系统基于Java和WML(无线标记语言)技术开发&#xff0c;实现了移动设备上的信息查询与后端信息发布功能。系统采用B/S架构&#xff0c;服务器端使用Java Servlet处理请求&#xff0c;数据库采用MySQL存储信息&#xff0…...

PHP 8.5 即将发布:管道操作符、强力调试

前不久&#xff0c;PHP宣布了即将在 2025 年 11 月 20 日 正式发布的 PHP 8.5&#xff01;作为 PHP 语言的又一次重要迭代&#xff0c;PHP 8.5 承诺带来一系列旨在提升代码可读性、健壮性以及开发者效率的改进。而更令人兴奋的是&#xff0c;借助强大的本地开发环境 ServBay&am…...

数学建模-滑翔伞伞翼面积的设计,运动状态计算和优化 !

我们考虑滑翔伞的伞翼面积设计问题以及运动状态描述。滑翔伞的性能主要取决于伞翼面积、气动特性以及飞行员的重量。我们的目标是建立数学模型来描述滑翔伞的运动状态,并优化伞翼面积的设计。 一、问题分析 滑翔伞在飞行过程中受到重力、升力和阻力的作用。升力和阻力与伞翼面…...

MySQL的pymysql操作

本章是MySQL的最后一章&#xff0c;MySQL到此完结&#xff0c;下一站Hadoop&#xff01;&#xff01;&#xff01; 这章很简单&#xff0c;完整代码在最后&#xff0c;详细讲解之前python课程里面也有&#xff0c;感兴趣的可以往前找一下 一、查询操作 我们需要打开pycharm …...