CUDA内存(三) 常量内存 __constant__

最新推荐文章于 2023-05-31 14:35:35 发布
最新推荐文章于 2023-05-31 14:35:35 发布
阅读量1.8k 收藏 2
点赞数 1
分类专栏： CUDA
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/full_speed_turbo/article/details/87173446
版权

文章目录

常量内存

常量内存其实只是全局内存的一种虚拟地址形式, 并没有特殊保留的常量内存块. 常量内存有两个特性, 一个是高速缓存, 另一个是它支持将单个值广播到线程束的每个线程.

测试代码

ConstMem.h

#pragma once

#include "Global.h"
#include <stdio.h>
#include <stdlib.h>
#include <conio.h>
#include <assert.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#define CUDA_CALL(x){const cudaError_t a=(x);if(a!=cudaSuccess){printf("\nCUDA Error: %s (err_num=%d)\n", cudaGetErrorString(a), a);cudaDeviceReset();assert(0);}} 

#define KERNEL_LOOP 65536

__constant__ static const u32 const_data_01 = 0x55555555;
__constant__ static const u32 const_data_02 = 0x77777777;
__constant__ static const u32 const_data_03 = 0x33333333;
__constant__ static const u32 const_data_04 = 0x11111111;

__device__ static  const u32 gmem_data_01 = 0x55555555;
__device__ static  const u32 gmem_data_02 = 0x77777777;
__device__ static  const u32 gmem_data_03 = 0x33333333;
__device__ static  const u32 gmem_data_04 = 0x11111111;

__global__ void const_test_gpu_literal(u32* const data, const u32 num_elements);
__global__ void const_test_gpu_const(u32 *const data, const u32 num_elements);
__global__ void const_test_gpu_gmem(u32 *const data, const u32 num_elements);
__host__ void wait_exit(void);
__host__ void cuda_error_check(const char *prefix, const char *postfix);
__host__ void gpu_const_vs_literal_kernel(void);
__host__ void gpu_const_vs_gmem_kernel(void);

ConstMem.cu


#include "ConstMem.h"

__global__ void const_test_gpu_literal(u32* const data, const u32 num_elements)
{
	const u32 tid = (blockIdx.x*blockDim.x) + threadIdx.x;
	if (tid < num_elements)
	{
		u32 d = 0x55555555;
		for (int i = 0; i < KERNEL_LOOP; i++)
		{
			d ^= 0x55555555;
			d |= 0x77777777;
			d &= 0x33333333;
			d |= 0x11111111;
		}
		data[tid] = d;
	}
}


__global__ void const_test_gpu_const(u32 *const data, const u32 num_elements)
{
	const u32 tid = (blockIdx.x*blockDim.x) + threadIdx.x;
	if (tid < num_elements)
	{
		u32 d = const_data_01;
		for (int i = 0; i < KERNEL_LOOP; i++)
		{
			d ^= const_data_01;
			d |= const_data_02;
			d &= const_data_03;
			d |= const_data_04;
		}
		data[tid] = d;
	}
}

__global__ void const_test_gpu_gmem(u32 *const data, const u32 num_elements)
{
	const u32 tid = (blockIdx.x*blockDim.x) + threadIdx.x;
	if (tid < num_elements)
	{
		u32 d = 0x55555555;
		for (int i = 0; i < KERNEL_LOOP; i++)
		{
			d ^= gmem_data_01;
			d |= gmem_data_02;
			d &= gmem_data_03;
			d |= gmem_data_04;
		}
		data[tid] = d;
	}
}
__host__ void wait_exit(void)
{
	char ch;
	printf("\nPress any key to exit");
	ch = getch();
}

__host__ void cuda_error_check(const char *prefix, const char *postfix)
{
	if (cudaPeekAtLastError() != cudaSuccess)
	{
		printf("\n%s%s%s", prefix, cudaGetErrorString(cudaGetLastError()), postfix);
		cudaDeviceReset();
		wait_exit();
		exit(1);
	}
}

__host__ void gpu_const_vs_literal_kernel(void)
{
	const u32 num_elements = 128 * 1024;
	const u32 num_threads = 256;
	const u32 num_blocks = (num_elements + num_threads - 1) / num_threads;
	const u32 num_bytes = num_elements * sizeof(u32);
	int max_device_num;
	const int max_runs = 6;
	
	CUDA_CALL(cudaGetDeviceCount(&max_device_num));
	for (int device_num = 0; device_num < max_device_num;device_num++)
	{
		CUDA_CALL(cudaSetDevice(device_num));
		for (int num_test = 0; num_test < max_runs; num_test++)
		{
			u32 *data_gpu = NULL;
			cudaEvent_t kernel_start1, kernel_stop1;
			cudaEvent_t kernel_start2, kernel_stop2;
			float delta_time1 = 0.0f, delta_time2 = 0.0f;
			struct cudaDeviceProp device_prop;
			char device_prefix[261];
			CUDA_CALL(cudaMalloc(&data_gpu, num_bytes));
			CUDA_CALL(cudaEventCreate(&kernel_start1));
			CUDA_CALL(cudaEventCreate(&kernel_start2));
			CUDA_CALL(cudaEventCreateWithFlags(&kernel_stop1, cudaEventBlockingSync));
			CUDA_CALL(cudaEventCreateWithFlags(&kernel_stop2, cudaEventBlockingSync));


			//printf("\nLauching %u blocks. %u threads", num_blocks, num_threads);
			CUDA_CALL(cudaGetDeviceProperties(&device_prop, device_num));
			sprintf(device_prefix, "ID:%d %s:", device_num, device_prop.name);

			// 字面值
			// Warm up run
			//printf("\nLauching literal kernel warm-up");
			const_test_gpu_literal<<<num_blocks, num_threads>>>(data_gpu, num_elements);
			cuda_error_check("Error ", " returned from literal startup kernel");
			//Do the literal kernel
			//printf("\nLauching literal kernel");
			CUDA_CALL(cudaEventRecord(kernel_start1, 0));
			const_test_gpu_literal << <num_blocks, num_threads >> >(data_gpu, num_elements);
			cuda_error_check("Error ", " returned from literal runtime kernel");
			CUDA_CALL(cudaEventRecord(kernel_stop1, 0));
			CUDA_CALL(cudaEventSynchronize(kernel_stop1));
			CUDA_CALL(cudaEventElapsedTime(&delta_time1, kernel_start1, kernel_stop1));
			//printf("\nLiteral Elapsed time: %.3fms", delta_time1);

			// 常量内存
			// Warm up run
			//printf("\nLauching constant kernel warm-up");
			const_test_gpu_const << <num_blocks, num_threads >> >(data_gpu, num_elements);
			cuda_error_check("Error ", " returned from costant startup kernel");

			// Do the const kernel
			//printf("\n Lauching constant kernel");
			CUDA_CALL(cudaEventRecord(kernel_start2, 0));
			const_test_gpu_const << <num_blocks, num_threads >> >(data_gpu, num_elements);
			cuda_error_check("Error ", " returned from constant runtime kernel");
			CUDA_CALL(cudaEventRecord(kernel_stop2, 0));
			CUDA_CALL(cudaEventSynchronize(kernel_stop2));
			CUDA_CALL(cudaEventElapsedTime(&delta_time2, kernel_start2, kernel_stop2));
			//printf("\nConstant Elapsed time: %.3fms", delta_time2);

			if (delta_time1 > delta_time2)
			{
				printf("\n%s常量内存快了: %.2fms (Const=%.2fms vs. Literal=%.2fms)", device_prefix, delta_time1 - delta_time2, delta_time1, delta_time2);
			}
			else
			{
				printf("\n%s字面取值快了: %.2fms (Const=%.2fms vs. Literal=%.2fms)", device_prefix, delta_time2 - delta_time1, delta_time1, delta_time2);
			}
			CUDA_CALL(cudaEventDestroy(kernel_start1));
			CUDA_CALL(cudaEventDestroy(kernel_start2));
			CUDA_CALL(cudaEventDestroy(kernel_stop1));
			CUDA_CALL(cudaEventDestroy(kernel_stop2));
			CUDA_CALL(cudaFree(data_gpu));

		}/* end of num_test*/
		CUDA_CALL(cudaDeviceReset());
		printf("\n");
	}/* end of device_num*/
}
__host__ void gpu_const_vs_gmem_kernel(void)
{
	const u32 num_elements = 128 * 1024;
	const u32 num_threads = 256;
	const u32 num_blocks = (num_elements + num_threads - 1) / num_threads;
	const u32 num_bytes = num_elements * sizeof(u32);
	int max_device_num;
	const int max_runs = 6;

	CUDA_CALL(cudaGetDeviceCount(&max_device_num));
	for (int device_num = 0; device_num < max_device_num; device_num++)
	{
		CUDA_CALL(cudaSetDevice(device_num));
		for (int num_test = 0; num_test < max_runs; num_test++)
		{
			u32 *data_gpu = NULL;
			cudaEvent_t kernel_start1, kernel_stop1;
			cudaEvent_t kernel_start2, kernel_stop2;
			float delta_time1 = 0.0f, delta_time2 = 0.0f;
			struct cudaDeviceProp device_prop;
			char device_prefix[261];
			CUDA_CALL(cudaMalloc(&data_gpu, num_bytes));
			CUDA_CALL(cudaEventCreate(&kernel_start1));
			CUDA_CALL(cudaEventCreate(&kernel_start2));
			CUDA_CALL(cudaEventCreateWithFlags(&kernel_stop1, cudaEventBlockingSync));
			CUDA_CALL(cudaEventCreateWithFlags(&kernel_stop2, cudaEventBlockingSync));


			//printf("\nLauching %u blocks. %u threads", num_blocks, num_threads);
			CUDA_CALL(cudaGetDeviceProperties(&device_prop, device_num));
			sprintf(device_prefix, "ID:%d %s:", device_num, device_prop.name);

			// 常量内存
			// Warm up run
			//printf("\nLauching constant kernel warm-up");
			const_test_gpu_const << <num_blocks, num_threads >> >(data_gpu, num_elements);
			cuda_error_check("Error ", " returned from costant startup kernel");

			// Do the const kernel
			//printf("\n Lauching constant kernel");
			CUDA_CALL(cudaEventRecord(kernel_start1, 0));
			const_test_gpu_const << <num_blocks, num_threads >> >(data_gpu, num_elements);
			cuda_error_check("Error ", " returned from constant runtime kernel");
			CUDA_CALL(cudaEventRecord(kernel_stop1, 0));
			CUDA_CALL(cudaEventSynchronize(kernel_stop1));
			CUDA_CALL(cudaEventElapsedTime(&delta_time1, kernel_start1, kernel_stop1));
			//printf("\nConstant Elapsed time: %.3fms", delta_time2);

			// 全局内存
			// Warm up run
			//printf("\nLauching gmem kernel warm-up");
			const_test_gpu_gmem << <num_blocks, num_threads >> >(data_gpu, num_elements);
			cuda_error_check("Error ", " returned from gmem startup kernel");
			//printf("\n Lauching gmem kernel");
			CUDA_CALL(cudaEventRecord(kernel_start2, 0));
			const_test_gpu_gmem << <num_blocks, num_threads >> >(data_gpu, num_elements);
			cuda_error_check("Error ", " returned from gmem runtime kernel");
			CUDA_CALL(cudaEventRecord(kernel_stop2, 0));
			CUDA_CALL(cudaEventSynchronize(kernel_stop2));
			CUDA_CALL(cudaEventElapsedTime(&delta_time2, kernel_start2, kernel_stop2));

			if (delta_time1 > delta_time2)
			{
				printf("\n%s常量内存快了: %.2fms (Const=%.2fms vs. Gmem=%.2fms)", device_prefix, delta_time1 - delta_time2, delta_time1, delta_time2);
			}
			else
			{
				printf("\n%s全局内存快了: %.2fms (Const=%.2fms vs. Gmem=%.2fms)", device_prefix, delta_time2 - delta_time1, delta_time1, delta_time2);
			}
			CUDA_CALL(cudaEventDestroy(kernel_start1));
			CUDA_CALL(cudaEventDestroy(kernel_start2));
			CUDA_CALL(cudaEventDestroy(kernel_stop1));
			CUDA_CALL(cudaEventDestroy(kernel_stop2));
			CUDA_CALL(cudaFree(data_gpu));

		}/* end of num_test*/
		CUDA_CALL(cudaDeviceReset());
		printf("\n");
	}/* end of device_num*/
}

ConstMem_main.cpp

#include "ConstMem.h"
int main()
{
	getCpuInfo();
	getGpuInfo();
	printf("\n常亮内存和字面值u32位操作速度对比:\n");
	gpu_const_vs_literal_kernel();
	gpu_const_vs_gmem_kernel();
	return 0;
}

测试结果

可以看出, (在较高版本架构上, )全局内存借助一级缓存也能达到与常量内存相同的访问速度!
在这里插入图片描述

Full_Speed_Turbo
关注
  • 1
    点赞
  • 踩
  • 2
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
专栏目录
[7] CUDA之常量内存与纹理内存
yohnyang的博客
05-26 542
[7] CUDA之常量内存与纹理内存
cuda_code.zip_CUDA求逆_GPU_cuda 矩阵_矩阵求逆cuda_逆矩阵
07-15
3. **内存层次**：CUDA有多种内存类型，包括全局内存（Global Memory）、共享内存（Shared Memory）、寄存器（Registers）和常量内存（Constant Memory）。根据数据访问模式选择合适的内存层次，可以显著提高性能。 ...
CUDA中的常量内存__constant__
热门推荐
牧野的博客
02-06 1万+
GPU包含数百个数学计算单元，具有强大的处理运算能力，可以强大到计算速率高于输入数据的速率，即充分利用带宽，满负荷向GPU传输数据还不够它计算的。CUDA C除全局内存和共享内存外，还支持常量内存，常量内存用于保存在核函数执行期间不会发生变化的数据，使用常量内存在一些情况下，能有效减少内存带宽，降低GPU运算单元的空闲等待。 使用常量内存提升性能 使用常量内存可以提升运算性能的
CUDA编程第五章: 共享内存&常量内存
Janus
02-16 5126
本章大概50页, 2天之内看完 前言: 本章内容： 了解数据在共享内存中是如何被安排的 掌握从二维共享内存到线性全局内存的索引转换 解决不同访问模式中存储体中的冲突 在共享内存中缓存数据以减少对全局内存的访问 使用共享内存避免非合并全局内存的访问 理解常量缓存和只读缓存之间的差异 使用线程束洗牌指令编程 在前面的章节中, 已经介绍了几种全局内存的访问模式. 通过安排全局内存访问模式, 我们学会了如何实现良好的性能并且避免了浪费事务. 未对齐的内存访问是没有问题的, 因为现.
cuda 学习之内存层次结构
xukang95的博客
03-23 2882
文章目录Cuda 内存层次结构registers：寄存器local memory：局部存储器shared memory：共享存储器共享内存-Bank Conflictconstant memory：常数存储器texture memory：纹理存储器global memory：全局存储器全局内存-对齐与合并host memory :主机端内存总结 Cuda 内存层次结构 要想编写高效的程序，那么一...
CUDA C编程（十五）常量内存
AAAA202012的博客
01-05 1921
   使 用 常 量 内 存 实 现 一 维 模 板    与 只 缓 内 存 的 比 较   
CUDA_常量内存
weixin_48151686的博客
11-09 725
常量内存简介 常量内存其实只是全局内存的一种虚拟地址形式，并没有特殊保留的常量内存块。常量内存有两个特性： 高速缓存,拥有缓存加速； 支持将单个值广播到线程束中的每个线程 常量内存的大小限定为64K，每个SM拥有8KB的常数存储器缓存，在编译期时声明一块常量内存，需要用到__constant__关键字，例如 __constant__ float my_array[1024] = { 0.0F, 1.0F, 1.3F, ...}; 不同于c/c++中的const常量，cuda中常量内存在声明后是可以修改
Cuda_SinglePage_site
08-04
- **常量内存（Constant Memory）**：存储不会改变的数据，访问速度快。 - **纹理内存（Texture Memory）**：优化了浮点数据的读取，适用于图像处理等场景。 CUDA编程还包括流（Streams）、事件（Events）、同步...
CUDA文档gpu应用.rar_GPU_gpu一下文档
09-21
7. **常量内存（Constant Memory）**：全局可读，写入后不会改变的内存区域。 8. **纹理内存（Texture Memory）**：优化了对2D数组访问的内存类型，适合于图像处理。 CUDA编程涉及的关键步骤： 1. **初始化CUDA**：...
Python 17_tensorflow.zip
02-23
例如，可以使用`tf.constant()`创建常量张量，`tf.Variable()`创建可训练的变量，以及使用`tf.add()`、`tf.subtract()`等函数进行加减运算。此外，张量还可以通过`tf.reshape()`改变形状，`tf.transpose()`转置，...
gpu_cpu_demo.tar.gz
06-21
3. **CUDA内存层次**：CUDA提供了多种内存类型，如全局内存（Global Memory）、共享内存（Shared Memory）、常量内存（Constant Memory）和纹理内存（Texture Memory），每种都有其特定的访问速度和使用场景。...
CUDA内存管理一文理清
HinGwenWoong的博客
02-13 1696
CUDA内存管理
CUDA进阶：内存模型
Shilong的博客
09-19 1171
CUDA内存模型
CUDA中的存储空间修饰符
专注于人工智能领域的小何尚
06-30 600
变量内存空间说明符表示变量在设备上的内存位置。在设备代码中声明的没有本节中描述的任何 、 和 内存空间说明符的自动变量通常驻留在寄存器中。 但是，在某些情况下，编译器可能会选择将其放置在本地内存中，这可能会产生不利的性能后果，如设备内存访问中所述。 内存空间说明符声明了一个驻留在设备上的变量。在接下来的三个部分中定义的其他内存空间说明符中最多有一个可以与 一起使用，以进一步表示变量属于哪个内存空间。 如果它们都不存在，则变量： 内存空间说明符，可选择与 一起使用，声明一个变量： 内存空间说明符，可选择
cuda编程学习——CUDA内存介绍（七）
出门吃三碗饭的博客
05-31 1298
现代计算机中的内存往往存在一种组织结构（hierarchy）。在这种结构中，含有多种类 型的内存，每种内存分别具有不同的容量和延迟（latency，可以理解为处理器等待内存数据的时间）。一般来说，延迟低（速度高）的内存容量小，延迟高（速度低）的内存容量大。
CUDA C编程（十一）CUDA共享内存概述
AAAA202012的博客
12-31 4555
  GPU有两种类型的内存：板载内存以及片上内存。其中全局内存是较大的板载内存，具有相对较高的延迟。共享内存是较小的片上内存，具有相对较低的延迟，并且共享内存可以提供比全局内存高得多的带宽。可以把它当作一个可编程管理的缓存，共享内存通常的用途有：块间线程通信的通道；用于全局内存数据的可编程管理的缓存；高速暂存存储器，用于转换数据以优化全局内存访问模式。 共 享 内 存    共 享 内 存 分 配 共 享 内 存 存 储 体 和 访 问 模 式 配 置 共 享 内 存 量 同 步 ...
CUDA——常量内存初始化
UCAS_HMM的博客
11-02 626
单变量常量内存CUDA，报错 invalid device symbol
CUDA Constant Memory常量内存使用问题
LuchangLi 的专栏
03-24 950
CUDA Constant Memory声明后，不能作为参数传递。 比如可以这样 __constant__ int constData[5]; __global__ void addKernel(int *c, const int *a, const int *b) { int i = threadIdx.x; c[i] = a[i] + b[i] + constData[i...
如何在CUDA中使用共享内存和常量内存？它们有什么作用和注意事项？
最新发布
06-12
使用共享内存和常量内存是在CUDA中优化内存访问的常用方法之一。 在CUDA中使用共享内存的方法如下： 1. 声明共享内存： ```__shared__ float shared_memory[SIZE];``` 2. 在内核中使用共享内存： ``` __global_...

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值