CUDA 学习笔记

最新推荐文章于 2021-03-29 11:26:14 发布

ipanda_zll

最新推荐文章于 2021-03-29 11:26:14 发布

阅读量508

点赞数

本文链接：https://blog.csdn.net/ipanda_huanhuan/article/details/107300804

版权

本文介绍了CUDA编程，包括在Windows10环境下使用VS2019和CUDA11.0进行编程的基础操作，如数据传输。重点讨论了统一内存的概念和使用，以及计算运行时间和错误检查。同时，深入探讨了CUDA的优化策略，如合并内存访问、减少线程束中的指令分歧、循环展开和使用cudaOccupancyMaxPotentialBlockSize计算最佳块大小。还提到了线程束、寄存器和共享内存在GPU架构中的重要性，强调了优化寄存器使用和利用共享内存提高性能的方法。

摘要由CSDN通过智能技术生成

windows10 VS2019 CUDA11.0
1.主流的标准操作：创建一些额外的数组并显式地使用cudaMemcpy（）函数，在主机和设备间通过PCIe总线进行数据的输入和输出。

//程序改编自书籍：《CUDA高性能并行计算》
#ifndef distance_h
#define distance_h

void distanceArray(float* out, float* in, float ref, int n);
float scale(int i, int n);

#endif

#include <stdio.h>
//也可不加此头文件，因为CUDA内联文件已经包含了math.h，程序可以正常运行
//只是不加此头文件，VS会因为识别不了sqrt()函数而报错
#include <cmath>
//也可不加以下头文件，程序可以正常运行，只是VS会报错
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#define TPB 32

//从核函数中调用并在GPU上执行，因此需要添加__device__标识。
__device__ float distance(float x1, float x2)//左右各两个下划线
{
   
	return sqrt((x2 - x1) * (x2 - x1));
}

//如所有的核函数一样，是从主机上调用而在设备上执行，因此需要使用修饰符__global__，并且设置返回类型为void。
__global__ void distanceKernal(float* d_out, float* d_in, float ref)
{
   
	const int i = blockIdx.x * blockDim.x + threadIdx.x;//计算出线程索引号，然后分散到各线程里进行并行计算
	const float x = d_in[i];
	d_out[i] = distance(x, ref);
	printf("i=%2d: dist from %f to %f is %f.\n", i, ref, x, d_out[i]);
}

//以下为主流的标准操作：创建一些额外的数组并显式地使用cudaMemcpy（）函数，在主机和设备间通过PCIe总线进行数据的输入和输出。

//函数distanceArray()调用核函数distanceKernel()。这样的函数被称为封装（warpper）或者启动函数（launcher）。
__host__ void distanceArray(float* out, float* in, float ref, int n)
{
   
	//声明指向设备端数组的指针
	float* d_in = NULL;
	float* d_out = NULL;
	//在设备端为设备端数组分配内存
	cudaMalloc(&d_in, n * sizeof(float));
	cudaMalloc(&d_out, n * sizeof(float));
	//把主机端上的数组复制到设备端对应的数组中
	cudaMemcpy(d_in, in, n * sizeof(float), cudaMemcpyHostToDevice