《基于 CUDA 的 GPU 并行程序开发指南》中的 Hflip7 的越界错误的原因

于 2023-12-26 10:44:19 发布
阅读量806 收藏 23
点赞数 15
分类专栏： CUDA 文章标签： c++ CUDA
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/hijackedbycsdn/article/details/135206465
版权

目录

书名：《基于 CUDA 的 GPU 并行程序开发指南》

越界问题

实际运行发现有越界的报错，经过测试，是 MYdstIndex 可能为 -1，转成 unsigned 变成了 2^32 - 1 ，是一个很大的数

计算出 MYdstIndex 判断一下他是不是负数就行了

// Improved Hflip6() kernel that flips the given image horizontally
// Each kernel: uses Shared Memory (PixBuffer[]) to read in 12 Bytes
// (4 pixels). 12Bytes are flipped inside Shared Memory 
// After that, they are written into Global Mem as 3 int's
// Horizontal resolution MUST BE A POWER OF 4.
__global__
void Hflip7(ui* ImgDst32, ui* ImgSrc32, ui RowInts)
{
	__shared__ ui PixBuffer[3072]; // holds 3*1024*4 Bytes (1024*4 pixels).

	ui ThrPerBlk = blockDim.x;
	ui MYbid = blockIdx.x;
	ui MYtid = threadIdx.x;
	ui MYtid3 = MYtid * 3;
	ui MYrow = blockIdx.y;
	ui MYcolIndex = (MYbid * ThrPerBlk + MYtid) * 3;
	if (MYcolIndex >= RowInts) return;			// index is out of range
	ui MYmirrorcol = RowInts - 1 - MYcolIndex;
	ui MYoffset = MYrow * RowInts;
	ui MYsrcIndex = MYoffset + MYcolIndex;
	ui MYdstIndex = MYoffset + MYmirrorcol - 2; // -2 is to copy 3 Bytes at a time
	if ((int)MYdstIndex < 0) return;

	uch SwapB;
	uch* SwapPtr;

	// read 4 pixel blocks (12B = 3 int's) into Shared Memory
	// PixBuffer:  [B0 G0 R0 B1] [G1 R1 B2 G2] [R2 B3 G3 R3]
	// Our Target: [B3 G3 R3 B2] [G2 R2 B1 G1] [R1 B0 G0 R0]
	PixBuffer[MYtid3] = ImgSrc32[MYsrcIndex];
	PixBuffer[MYtid3 + 1] = ImgSrc32[MYsrcIndex + 1];
	PixBuffer[MYtid3 + 2] = ImgSrc32[MYsrcIndex + 2];
	__syncthreads();

	// swap these 4 pixels inside Shared Memory
	SwapPtr = (uch*)(&PixBuffer[MYtid3]);      // [B0 G0 R0 B1] [G1 R1 B2 G2] [R2 B3 G3 R3]
	SWAP(SwapPtr[0], SwapPtr[9], SwapB)			// [B3 G0 R0 B1] [G1 R1 B2 G2] [R2 B0 G3 R3]
	SWAP(SwapPtr[1], SwapPtr[10], SwapB)		// [B3 G3 R0 B1] [G1 R1 B2 G2] [R2 B0 G0 R3]
	SWAP(SwapPtr[2], SwapPtr[11], SwapB)		// [B3 G3 R3 B1] [G1 R1 B2 G2] [R2 B0 G0 R0]
	SWAP(SwapPtr[3], SwapPtr[6], SwapB)			// [B3 G3 R3 B2] [G1 R1 B1 G2] [R2 B0 G0 R0]
	SWAP(SwapPtr[4], SwapPtr[7], SwapB)			// [B3 G3 R3 B2] [G2 R1 B1 G1] [R2 B0 G0 R0]
	SWAP(SwapPtr[5], SwapPtr[8], SwapB)			// [B3 G3 R3 B2] [G2 R2 B1 G1] [R1 B0 G0 R0]

	__syncthreads();

	//write the 4 pixels (3 int's) from Shared Memory into Global Memory
	ImgDst32[MYdstIndex] = PixBuffer[MYtid3];
	ImgDst32[MYdstIndex + 1] = PixBuffer[MYtid3 + 1];
	ImgDst32[MYdstIndex + 2] = PixBuffer[MYtid3 + 2];
}

序号计算逻辑

其实虽然我发现了是哪个序号出现了负数，但是我不知道他为什么会是负的

所以这里做一个分析

最初的函数例子

#define	IPH			ip.Hpixels
#define	IPV			ip.Vpixels
#define	CEIL(a,b)		((a+b-1)/b)

...

BlkPerRow = CEIL(IPH,ThrPerBlk);
NumBlocks = IPV*BlkPerRow;

...

Hflip <<< NumBlocks, ThrPerBlk >>> (GPUCopyImg, GPUImg, IPH);

他这里是一个线程处理一个像素的，所以基本思路就是你至少要启动跟像素总数相等的线程

然后有又因为线程块里面的线程数是固定的，所以肯定会有浪费

所以要向上取整

现在把他拆成一个二维的，对第一个维度除以 4，那就是一个线程处理 4 个像素

dim3 dimGrid2D4(CEIL(BlkPerRow,4), ip.Vpixels);

如果是一个线程处理一个像素

ui ThrPerBlk = blockDim.x;
ui MYbid = blockIdx.x;
ui MYtid = threadIdx.x;

ui MYrow = blockIdx.y;
ui MYcol = MYbid*ThrPerBlk + MYtid;

这里算出来的 MYcol 就是当前行要处理的像素的列数

但是一个线程处理 4 个像素，一个像素 3 byte，4 个像素就是 12 byte，也等于 3 个 int

那么现在我不算当前行要处理的像素的列数，我把当前行的像素总数换算成 int 的个数

然后我现在要算的是当前行要处理的 int 的列数，也就是要处理第几个 int

ui ThrPerBlk = blockDim.x;
ui MYbid = blockIdx.x;
ui MYtid = threadIdx.x;

ui MYrow = blockIdx.y;
ui MYcolIndex = (MYbid * ThrPerBlk + MYtid) * 3;

这个 MYcolIndex 就是当前行要处理的 int 的序号

那么接下来就很清楚了

ui MYoffset = MYrow * RowInts;
ui MYsrcIndex = MYoffset + MYcolIndex;
ui MYdstIndex = MYoffset + MYmirrorcol - 2;

当前行要处理的 int 的序号显然可以是 0 或者 1，如果是第 0 行的话，那么 MYoffset = 0

那么显然 MYdstIndex 可能是 -1 或者 -2

所以要加这个判断条件

水平翻转上的错误

原图

在这里插入图片描述

水平翻转之后的结果

在这里插入图片描述

这种错误只有我在开始用 flip7 之后才出现的

所以还是需要思考一下为什么会这样……

想了一天似乎还是没搞懂

比如我把程序开始和结束时的 img 指针的内容打印出来

...

printf("TheImg:\n");
for (int i = 0; i < IPV; ++i) {
	for (int j = 0; j < IPH; ++j) {
		printf("(%d, ", *(TheImg + i * IPHB + 3 * j));
		printf("%d, ", *(TheImg + i * IPHB + 3 * j + 1));
		printf("%d) ", *(TheImg + i * IPHB + 3 * j + 2));
	}
	printf("\n");
}

...

printf("CopyImg:\n");
for (int i = 0; i < IPV; ++i) {
	for (int j = 0; j < IPH; ++j) {
		printf("(%d, ", *(CopyImg + i * IPHB + 3 * j));
		printf("%d, ", *(CopyImg + i * IPHB + 3 * j + 1));
		printf("%d) ", *(CopyImg + i * IPHB + 3 * j + 2));
	}
	printf("\n");
}

输出：

TheImg:
(255, 255, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255) (0, 0, 255)
(255, 255, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255) (0, 0, 255) (255, 0, 0)
(255, 255, 255) (255, 255, 255) (255, 255, 255) (0, 0, 255) (255, 0, 0) (0, 255, 0)
(255, 255, 255) (255, 255, 255) (0, 0, 255) (255, 0, 0) (0, 255, 0) (255, 255, 255)
(255, 255, 255) (0, 0, 255) (255, 0, 0) (0, 255, 0) (255, 255, 255) (255, 255, 255)
(0, 0, 255) (255, 0, 0) (0, 255, 0) (255, 255, 255) (255, 255, 255) (255, 255, 255)
CopyImg:
(0, 0, 0) (0, 0, 0) (0, 0, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255)
(0, 255, 255) (0, 0, 0) (0, 255, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255)
(0, 255, 0) (255, 0, 255) (0, 0, 0) (0, 255, 255) (255, 255, 255) (255, 255, 255)
(0, 255, 255) (255, 255, 0) (255, 0, 255) (0, 0, 0) (0, 255, 255) (255, 255, 255)
(0, 0, 255) (255, 255, 255) (255, 255, 0) (255, 0, 255) (0, 0, 0) (0, 255, 255)
(0, 0, 255) (255, 255, 255) (255, 255, 255) (255, 255, 0) (255, 0, 255) (0, 0, 0)

这个输出的行的方向是从图的底部到图的顶部

如果在 kernel 里面加一个判断，只允许第 0 行和第 0 列的执行

if (MYsrcIndex != 0 || MYrow != 0) return;

TheImg:
(255, 255, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255) (0, 0, 255)
(255, 255, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255) (0, 0, 255) (255, 0, 0)
(255, 255, 255) (255, 255, 255) (255, 255, 255) (0, 0, 255) (255, 0, 0) (0, 255, 0)
(255, 255, 255) (255, 255, 255) (0, 0, 255) (255, 0, 0) (0, 255, 0) (255, 255, 255)
(255, 255, 255) (0, 0, 255) (255, 0, 0) (0, 255, 0) (255, 255, 255) (255, 255, 255)
(0, 0, 255) (255, 0, 0) (0, 255, 0) (255, 255, 255) (255, 255, 255) (255, 255, 255)
CopyImg:
(0, 0, 0) (0, 0, 0) (0, 0, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255)
(0, 0, 0) (0, 0, 0) (0, 0, 0) (0, 0, 0) (0, 0, 0) (0, 0, 0)
(0, 0, 0) (0, 0, 0) (0, 0, 0) (0, 0, 0) (0, 0, 0) (0, 0, 0)
(0, 0, 0) (0, 0, 0) (0, 0, 0) (0, 0, 0) (0, 0, 0) (0, 0, 0)
(0, 0, 0) (0, 0, 0) (0, 0, 0) (0, 0, 0) (0, 0, 0) (0, 0, 0)
(0, 0, 0) (0, 0, 0) (0, 0, 0) (0, 0, 0) (0, 0, 0) (0, 0, 0)

如果是改成

if (MYcolIndex != 0) return;

输出 2

TheImg:
(255, 255, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255) (0, 0, 255)
(255, 255, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255) (0, 0, 255) (255, 0, 0)
(255, 255, 255) (255, 255, 255) (255, 255, 255) (0, 0, 255) (255, 0, 0) (0, 255, 0)
(255, 255, 255) (255, 255, 255) (0, 0, 255) (255, 0, 0) (0, 255, 0) (255, 255, 255)
(255, 255, 255) (0, 0, 255) (255, 0, 0) (0, 255, 0) (255, 255, 255) (255, 255, 255)
(0, 0, 255) (255, 0, 0) (0, 255, 0) (255, 255, 255) (255, 255, 255) (255, 255, 255)
CopyImg:
(0, 0, 0) (0, 0, 0) (0, 0, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255)
(0, 0, 0) (0, 0, 0) (0, 0, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255)
(0, 0, 0) (0, 0, 0) (0, 0, 0) (0, 255, 255) (255, 255, 255) (255, 255, 255)
(0, 0, 0) (0, 0, 0) (0, 0, 255) (0, 0, 0) (0, 255, 255) (255, 255, 255)
(0, 0, 0) (0, 0, 0) (0, 0, 0) (255, 0, 255) (0, 0, 0) (0, 255, 255)
(0, 0, 0) (0, 0, 0) (0, 0, 255) (255, 255, 0) (255, 0, 255) (0, 0, 0)

可见他这里是，每行的后面被吞了 2 个字节

把最终的输出排一下版

(0, 0, 0) (0, 0, 0)         (0, 0, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255)
(0, 255, 255) (0, 0, 0)     (0, 255, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255)
(0, 255, 0) (255, 0, 255)   (0, 0, 0) (0, 255, 255) (255, 255, 255) (255, 255, 255)
(0, 255, 255) (255, 255, 0) (255, 0, 255) (0, 0, 0) (0, 255, 255) (255, 255, 255)
(0, 0, 255) (255, 255, 255) (255, 255, 0) (255, 0, 255) (0, 0, 0) (0, 255, 255)
(0, 0, 255) (255, 255, 255) (255, 255, 255) (255, 255, 0) (255, 0, 255) (0, 0, 0)

起码确认了从 输出 2 到最终的输出，输出 2 这里的趋于没有被破坏

那就把最后面的也输出

...

printf("TheImg:\n");
for (int i = 0; i < IPV; ++i) {
	for (int j = 0; j < IPH; ++j) {
		printf("(%d, ", *(TheImg + i * IPHB + 3 * j));
		printf("%d, ", *(TheImg + i * IPHB + 3 * j + 1));
		printf("%d) ", *(TheImg + i * IPHB + 3 * j + 2));
	}

	for (int j = 3 * IPH; j < IPHB; ++j) {
		printf("%d ", *(TheImg + i * IPHB + j));
	}
	printf("\n");
}

...

printf("CopyImg:\n");
for (int i = 0; i < IPV; ++i) {
	for (int j = 0; j < IPH; ++j) {
		printf("(%d, ", *(CopyImg + i * IPHB + 3 * j));
		printf("%d, ", *(CopyImg + i * IPHB + 3 * j + 1));
		printf("%d) ", *(CopyImg + i * IPHB + 3 * j + 2));
	}

	for (int j = 3 * IPH; j < IPHB; ++j) {
		printf("%d ", *(CopyImg + i * IPHB + j));
	}
	printf("\n");
}

输出

TheImg:
(255, 255, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255) (0, 0, 255) 0 0
(255, 255, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255) (0, 0, 255) (255, 0, 0) 0 0
(255, 255, 255) (255, 255, 255) (255, 255, 255) (0, 0, 255) (255, 0, 0) (0, 255, 0) 0 0
(255, 255, 255) (255, 255, 255) (0, 0, 255) (255, 0, 0) (0, 255, 0) (255, 255, 255) 0 0
(255, 255, 255) (0, 0, 255) (255, 0, 0) (0, 255, 0) (255, 255, 255) (255, 255, 255) 0 0
(0, 0, 255) (255, 0, 0) (0, 255, 0) (255, 255, 255) (255, 255, 255) (255, 255, 255) 0 0
CopyImg:
(0, 0, 0) (0, 0, 0) (0, 0, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255) 255 255
(0, 0, 0) (0, 0, 0) (0, 0, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255) 255 255
(0, 0, 0) (0, 0, 0) (0, 0, 0) (0, 255, 255) (255, 255, 255) (255, 255, 255) 255 255
(0, 0, 0) (0, 0, 0) (0, 0, 255) (0, 0, 0) (0, 255, 255) (255, 255, 255) 255 255
(0, 0, 0) (0, 0, 0) (0, 0, 0) (255, 0, 255) (0, 0, 0) (0, 255, 255) 255 255
(0, 0, 0) (0, 0, 0) (0, 0, 255) (255, 255, 0) (255, 0, 255) (0, 0, 0) 0 255

这就证明了确实是每行的后面被吞了 2 个字节

假如把核函数中的条件改成

if (MYcolIndex == 0) return;

输出：

TheImg:
(255, 255, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255) (0, 0, 255) 0 0
(255, 255, 255) (255, 255, 255) (255, 255, 255) (255, 255, 255) (0, 0, 255) (255, 0, 0) 0 0
(255, 255, 255) (255, 255, 255) (255, 255, 255) (0, 0, 255) (255, 0, 0) (0, 255, 0) 0 0
(255, 255, 255) (255, 255, 255) (0, 0, 255) (255, 0, 0) (0, 255, 0) (255, 255, 255) 0 0
(255, 255, 255) (0, 0, 255) (255, 0, 0) (0, 255, 0) (255, 255, 255) (255, 255, 255) 0 0
(0, 0, 255) (255, 0, 0) (0, 255, 0) (255, 255, 255) (255, 255, 255) (255, 255, 255) 0 0
CopyImg:
(0, 0, 0) (0, 0, 0) (0, 0,         0) (0, 0, 0) (0, 0, 0) (0, 255, 255) 255 0
(0, 255, 255) (0, 0, 0) (0, 255,   0) (0, 0, 0) (0, 0, 0) (0, 255, 255) 255 0
(0, 255, 0) (255, 0, 255) (0, 0,     0) (0, 0, 0) (0, 0, 0) (0, 255, 255) 0 0
(0, 255, 255) (255, 255, 0) (255, 0, 0) (0, 0, 0) (0, 0, 0) (0, 0, 255) 255 0
(0, 0, 255) (255, 255, 255) (255, 255,   0) (0, 0, 0) (0, 0, 0) (0, 0, 0) 0 0
(0, 0, 255) (255, 255, 255) (255, 255,   0) (0, 0, 0) (0, 0, 0) (0, 0, 0) 0 0

这就是为什么明明 MYmirrorcol - 2 会为负，但是只有一个位置会报错，

因为算

ui MYsrcIndex = MYoffset + MYcolIndex;
ui MYdstIndex = MYoffset + MYmirrorcol - 2;

的时候还加上了一个 MYoffset

这样的话，其实如果你的 MYmirrorcol - 2 会为负数的话，其实会发生一件事就是，MYsrcIndex 涉及到的范围跨越了两行，MYdstIndex 涉及到的范围也跨越了两行

比如这个输出中可以看出

在这里插入图片描述
在这里插入图片描述

这样就……完全错了

现在我觉得我完全知道他为什么错了……但是我不想改了，已经花了很多时间了……

hijackedbycsdn
关注
专栏目录
《基于CUDA的GPU并行程序开发指南》阅读笔记
liferecords的博客
10-27 1196
《基于CUDA的GPU并行程序开发指南》旨在帮助读者了解与基于CUDA的GPU并行编程技术有关的基本概念，并掌握使用C语言进行GPU高性能编程的相关技巧。
CUDA学习--矩阵乘法的并行运算
热门推荐
Yi_M的博客
08-24 1万+
1. CUDA学习步骤 CPU实现 a*b = c 的矩阵乘法（矩阵尺寸是n*m的，n和m大于1000） 下载 https://developer.nvidia.com/cuda-downloads，安装好cuda 将cpu代码移植到cuda。将CPU值传入GPU，使用cuda计算，与cpu结果对比。 优化思路1：将矩阵分块进行计算 优化思路2：使用share memory进行优化 优化思路3：将数
如何处理GPU上Error Number:700 an illegal memory access was encounter
Kenji_Shinji的博客
09-13 2692
步骤2：执行步骤1的操作后问题不复现，也就是同步执行的这种方式下没有问题，基本上可以确定是有算子里依赖同步执行的结果，因为正常device算子执行是异步执行，所以拿的结果不是预期值，同步执行就掩盖了这个问题。步骤1：由于GPU的算子执行是host下发到device上异步执行的，因此执行报错的地方不一定是真凶，大概率是前面的算子有问题，但是device是异步执行的，所以执行到后面才会报错。案例代码：https://gitee.com/mindspore/mindspore/pulls/962。
《基于 CUDA 的 GPU 并行程序开发指南》中的 RowBytes 为什么这么计算
最新发布
hijackedbycsdn的博客
12-24 394
为什么要取 4 的倍数，这不是 CUDA 的要求，只是因为 bmp 的文件格式要求每行是 4 的整数倍。虽然你看到的文件属性显示的行像素值不是 4 的倍数，但是实际上文件的内容是把行填充到 4 的倍数了。书名：《基于 CUDA 的 GPU 并行程序开发指南》最简单的例子就是，如果你要读取，那么你是一次性读。相当于 A 向上取整取 4 的倍数。也就是取 4 的倍数。
基于C++与CUDA的N卡GPU并行程序——在python中使用numba库编写GPU程序
wanchaochaochao的博客
06-19 1313
文档   numba文档   nvidia论坛文档   numba官网简介 解决numba报错   在python中使用numba编写CUDA程序时会有一个报错 NvvmSupportError: libNVVM cannot be found. Do conda install cudatoolkit: 这一般是关于CUDA的环境变量没有识别出来,所以需要在bashrc或/etc/profile中加入环境变量 export CUDA_HOME=/usr/local/cuda expor
CUDA并行程序设计 GPU编程指南: 第一章：超级计算简史
qq_42738654的博客
10-27 463
超级计算机简史
CUDA并行性能分析与调试技巧指南
CUDA（Compute Unified Device Architecture）是英伟达推出的一种并行计算平台和编程模型，旨在利用GPU的并行计算能力来加速通用目的计算。本章将介绍CUDA的基础知识，包括其简介与发展历程、并行计算模型概述以及...
CUDA中的并行计算容器与容器优化
CUDA（Compute Unified Device Architecture）是由NVIDIA推出的一种并行计算平台和编程模型，它使得在NVIDIA的GPU（Graphics Processing Unit）上进行并行计算变得更加简单和高效。相比传统的CPU，在GPU上进行并行...
基于cuda的积分图均值滤波
qq_43339160的博客
09-30 1633
基于cuda的积分图均值滤波均值滤波原理积分图CPU版本一CPU版本二GPU版本版本比较算法分析 均值滤波原理 在卷积形式的均值滤波中，需要把输入图对应模板内所有点像素点相加求平均，代替原像素点。 对于输入输出图像大小相同的均值滤波算法而言。边界的处理情况是不同的。超出图像范围的部分进行补零，这会使边界的像素值在均值滤波之后均值变小。以图一为例，左上角四个元素的均值为4。边界补零后计算3×3区域内的均值为1.8，明显均值偏小。本文算法边界采用自适应方式，比如该例子只计算四个像素的均值。 积分图 图像积分图
NVIDIA-CUDA HPC 编程模型与内存管理初探
tugouxp的专栏
11-12 3171
主存和现存在PCIE框架下可以做到互相访问，路径归纳如下：三种角色，两种存储，六种路径。1.HBM通过BAR透给了HOST，使HOST可以通过MMU给CPU访问，或者通过IOMMU给设备访问。2.相反，Host Memory并没有类似的BAR机制透给 PCIE设备端，PCIE设备端想要访问 HOST MEMORY必须经过 HOST IOMMU的映射一遍，加上PCIE端的SMMU映射，对于PCIE设备端（GPU）来说，要经过两层IOMMU翻译。所以从HOST端和GPU端看起来不太对称。
基于GPU的并行CUDA编程
12-29
GPU并行计算原理，包括环境搭建，CUDA库的讲解，文件中有相关的学习代码。
CUDA并行程序设计 GPU编程指南 pdf 中文版 完整版.part1
03-31
CUDA开发者社区技术总监亲自撰写，英伟达中国首批CUDA官方认证工程师翻译，译著双馨 全面、详实地讲解了CUDA并行程序设计的技术知识点和编程方法，包含大量实用代码示例，是目前学习CUDA编程最权威的著作之一
GPU编程学习入门
千人斩的博客
10-05 5623
本文使用的显卡是Nvidia,编程语言是CUDA。 本文主要是记录一下学习GPU编程的过程，在一到两天的时间里，对GPU编程思想有一个大概的认识，为将来深入学习打下基础，也对GPU编程可以解决的问题边界有一个认知，不被讲PPT的忽悠。
NVIDIA GPU Display Driver 越界写入漏洞（CVE-2022-28181）
murphysec的博客
05-18 777
CVE-2022-28181 漏洞是 NVIDIA GPU Display Driver内核模式层中的漏洞，导致网络上的非特权普通用户可以通过crafted shader越界写入。该漏洞影响范围一般，但影响的范围可能会扩展到其他组件。官方已发布新版本，建议受影响用户升级到安全版本。 编号      CVE-2022-28181 标题 NVIDIA GPU Display Driver 越界写入漏洞 描述 英伟达（NVIDIA）是一家专注于交互式图形产品的公司。其适用于 Windows 和
对tensor不同的索引方法和索引越界问题解决思路
lyp82nlf666的博客
04-27 746
对tensor不同的索引方法和索引越界问题解决思路
CUDA中内存访问越界问题
fb_help的专栏
03-18 4927
CUDA中内存访问越界问题 在核函数中对非法的内存访问不一定会报错，那我们加CHECK（error）会帮我们报错吗？经实验发现，有的非法访问很致命，函数不会运行，这种错误，可以CHECK到，但有些非法内存访问不致命，可能是访问到其他设备内存上了，这时CHECK不会报错，但是会出现逻辑上的错误，因为访问的值是错误的。 所以，我们在做内存访问时，一定要自己做好逻辑上的判断，保证内存访问不错位，不错...
基于cuda的gpu并行程序开发指南 pdf_Python之父推荐，这本书籍堪称python入门宝藏！赠PDF电子版...
weixin_39631445的博客
12-07 664
人生苦短，我用Python。今天推荐的这本书，连python之父都说它好，认为它确实是值得一读的Python书籍。此书在简介中说明，阅读本书不需要任何的互联网协议知识。这本书的名字叫《Head First Python》（重视大脑的学习指南），是一位名叫：Paul Barry的大牛所编著的。最近发现他的编程生涯已近四分之一个世纪，这个事实着实让人有些震惊。在此期间，Paul使用过多种不同的编程语言...
基于cuda的gpu并行程序开发指南 pdf
12-07
《基于CUDA的GPU并行程序开发指南》是一本专门介绍使用CUDA技术来开发并行程序的指南。CUDA是一种由英伟达推出的并行计算平台和编程模型，它可以让开发者利用GPU的并行计算能力来加速程序运行速度。这本指南主要介绍...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值