1.写在前面
因为自己需要实现cuda编程做中心线的提取，实现灰度重心法的实验。所以学习cuda，但是本人从未接触过cuda，所以在这里记录下来学习cuda的过程。
2024.11.26，先学习一篇cuda编程入门极简教程，在csdn搜索一篇教程，先尝试跟着走一遍。教程如下：CUDA编程入门极简教程_cuda编程举例-优快云博客
学习教程画重点：
2.初识cuda
GPU并不是一个独立运行的计算平台，而需要与CPU协同工作，可以看成是CPU的协处理器，因此当我们在说GPU并行计算时，其实是指的基于CPU+GPU的异构计算架构。CPU所在位置称为为主机端（host），而GPU所在位置称为设备端（device）。
基于CPU+GPU的异构计算平台可以优势互补，CPU负责处理逻辑复杂的串行程序，而GPU重点处理数据密集型的并行计算程序。
典型的CUDA程序的执行流程如下：
- 分配host内存，并进行数据初始化；
- 分配device内存，并从host将数据拷贝到device上；
- 调用CUDA的核函数在device上完成指定的运算；
- 将device上的运算结果拷贝到host上；
- 释放device和host上分配的内存。
上面流程中最重要的一个过程是调用CUDA的核函数来执行并行计算，kernel是CUDA中一个重要的概念，kernel是在device上线程中并行执行的函数，核函数用__global__符号声明，在调用时需要用<<<grid, block>>>来指定kernel要执行的线程数量，在CUDA中，每一个线程都要执行核函数，并且每个线程会分配一个唯一的线程号thread ID，这个ID值可以通过核函数的内置变量threadIdx来获得。
__global__：在device上执行，从host中调用（一些特定的GPU也可以从device上调用），返回类型必须是void，不支持可变参数参数，不能成为类成员函数。注意用__global__定义的kernel是异步的，这意味着host不会等待kernel执行完就执行下一步。
__device__：在device上执行，单仅可以从device中调用，不可以和__global__同时用。
__host__：在host上执行，仅可以从host上调用，一般省略不写，不可以和__global__同时用，但可和__device__，此时函数会在device和host都编译。
想要根据博主的代码进行测试，但是发现自己还不知道 在哪里进行cuda编程。于是在vs中新建了一个cuda的源文件，文件后缀名字为*.cu
然后需要配置cuda的环境，因为cuda在安装时就已经自动的加入了系统环境变量，所以这里只需要添加属性配置。
配置如下：
1.包含目录：C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.0\include
库目录：C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.0\lib\x64
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.0\lib
2.附加依赖项：
cublas.lib
cublasLt.lib
cuda.lib
cudadevrt.lib
cudart.lib
cudart_static.lib
cufft.lib
cufftw.lib
curand.lib
cusolver.lib
cusolverMg.lib
cusparse.lib
nppc.lib
nppial.lib
nppicc.lib
nppidei.lib
nppif.lib
nppig.lib
nppim.lib
nppist.lib
nppisu.lib
nppitc.lib
npps.lib
nvblas.lib
nvjpeg.lib
nvml.lib
nvrtc.lib
OpenCL.lib
然后使用测试代码：代码出处：VS2019+CUDA编程（流程）_vs如何运行cu文件-优快云博客
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <stdlib.h>
void main() {
int deviceCount;
cudaGetDeviceCount(&deviceCount);
int dev;
for (dev = 0; dev < deviceCount; dev++)
{
int driver_version(0), runtime_version(0);
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, dev);
if (dev == 0)
if (deviceProp.minor = 9999 && deviceProp.major == 9999)
printf("\n");
printf("\nDevice%d:\"%s\"\n", dev, deviceProp.name);
cudaDriverGetVersion(&driver_version);
printf("CUDA驱动版本: %d.%d\n", driver_version / 1000, (driver_version % 1000) / 10);
cudaRuntimeGetVersion(&runtime_version);
printf("CUDA运行时版本: %d.%d\n", runtime_version / 1000, (runtime_version % 1000) / 10);
printf("设备计算能力: %d.%d\n", deviceProp.major, deviceProp.minor);
printf("Total amount of Global Memory: %u bytes\n", deviceProp.totalGlobalMem);
printf("Number of SMs: %d\n", deviceProp.multiProcessorCount);
printf("Total amount of Constant Memory: %u bytes\n", deviceProp.totalConstMem);
printf("Total amount of Shared Memory per block: %u bytes\n", deviceProp.sharedMemPerBlock);
printf("Total number of registers available per block: %d\n", deviceProp.regsPerBlock);
printf("Warp size: %d\n", deviceProp.warpSize);
printf("Maximum number of threads per SM: %d\n", deviceProp.maxThreadsPerMultiProcessor);
printf("Maximum number of threads per block: %d\n", deviceProp.maxThreadsPerBlock);
printf("Maximum size of each dimension of a block: %d x %d x %d\n", deviceProp.maxThreadsDim[0],
deviceProp.maxThreadsDim[1],
deviceProp.maxThreadsDim[2]);
printf("Maximum size of each dimension of a grid: %d x %d x %d\n", deviceProp.maxGridSize[0], deviceProp.maxGridSize[1], deviceProp.maxGridSize[2]);
printf("Maximum memory pitch: %u bytes\n", deviceProp.memPitch);
printf("Texture alignmemt: %u bytes\n", deviceProp.texturePitchAlignment);
printf("Clock rate: %.2f GHz\n", deviceProp.clockRate * 1e-6f);
printf("Memory Clock rate: %.0f MHz\n", deviceProp.memoryClockRate * 1e-3f);
printf("Memory Bus Width: %d-bit\n", deviceProp.memoryBusWidth);
}
system("pause");
//return 0;
}最终我的运行结果：
到此步骤，顺利输出显卡的参数信息，则证明环境没有问题，并且可以顺利的视线编译。
下一步继续学习如何使用cuda开辟空间和编程。
2024.11.29继续学习
已经学习了如何在vs中实现cuda文件如何实现函数进行编程，下面进行一个简单的实战，利用cuda实现两个向量的加法。
在实现之前，cuda最重要的内存管理，首先是在device上进行内存分配，使用的函数是cudaMalloc
cudaError_t cudaMalloc(void** devPtr, size_t size);
这个函数和C语言中的malloc类似，但是在device上申请一定字节大小的显存，其中devPtr是指向所分配内存的指针。同时要释放分配的内存使用cudaFree函数，这和C语言中的free函数对应。另外一个重要的函数是负责host和device之间数据通信的cudaMemcpy函数：
cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind)
其中src指向数据源，而dst是目标区域，count是复制的字节数，其中kind控制复制的方向：cudaMemcpyHostToHost, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost及cudaMemcpyDeviceToDevice，如cudaMemcpyHostToDevice将host上数据拷贝到device上。
其中src指向数据源，而dst是目标区域，count是复制的字节数，其中kind控制复制的方向：cudaMemcpyHostToHost, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost及cudaMemcpyDeviceToDevice，如cudaMemcpyHostToDevice将host上数据拷贝到device上。
加法实例
现在我们来实现一个向量加法的实例，这里grid和block都设计为1-dim，首先定义kernel如下：
// 两个向量加法kernel，grid和block均为一维
__global__ void add(float* x, float * y, float* z, int n)
{
// 获取全局索引
int index = threadIdx.x + blockIdx.x * blockDim.x;
// 步长
int stride = blockDim.x * gridDim.x;
for (int i = index; i < n; i += stride)
{
z[i] = x[i] + y[i];
}
}
其中stride是整个grid的线程数，有时候向量的元素数很多，这时候可以将在每个线程实现多个元素（元素总数/线程总数）的加法，相当于使用了多个grid来处理，这是一种grid-stride loop方式，不过下面的例子一个线程只处理一个元素，所以kernel里面的循环是不执行的。下面我们具体实现向量加法：
int main(){
int N = 1 << 20;//N被设置为2的20次方，等于1048576
int nBytes = N * sizeof(float);//计算所需内存字节数，float类型占用4字节。
//申请主机内存（Host Memory）
float *x, *y, *z;//申请3个浮点数数组的指针。
x = (float*)malloc(nBytes);//为x数组分配内存。
y= (float*)malloc(nBytes);//为y数组分配内存。
z = (float*)malloc(nBytes);//为z数组分配内存。
//初始化数据
for (int i=0;i<N;++i)
{
x[i] = 10.0;//将数组x的每个袁术初始化为10.0
y[i] = 20.0;//将数组x的每个袁术初始化为20.0
}
//申请设备内存（Device Memory）
float *d_x, *d_y, *d_z;//申明设备端浮点数数组的指针
cudaMalloc((void**)&d_x, nBytes);//在设备上为d_x分配内存
cudaMalloc((void**)&d_y, nBytes);//在设备上为d_y分配内存
cudaMalloc((void**)&d_z, nBytes);//在设备上为d_z分配内存
//将主机数据拷贝到设备：
cudaMemcpy((void*)d_x, (void*)x, nBytes, cudaMemcpyHostToDevice);//将x的数据从主机复制到设备。
cudaMemcpy((void*)d_y,(void*)y, nBytes,cudaMemcpyHostToDevice);//将y的数据从主机复制到设备。
//定义内核执行配置：
dim3 blockSize(256);//每个块的线程数量设置为256。
dim3 gridSize((N+ blockSize.x-1)/ blockSize.x);//计算网格大小，根据线程块的大小确定需要多少个块
//执行内核
add <<<gridSize, blockSize >>> (d_x, d_y, d_z,N);
//将设备计算结果拷贝回主机：
cudaMemcpy((void*)z,(void*)d_z, nBytes,cudaMemcpyDeviceToHost);//将d_z的数据拷贝回z。
//执行检查结果：
float maxError = 0.0;
for (int i = 0; i < nBytes; i++)
{
maxError = fmax(maxError,fabs(z[i]-30.0));//计算结果中最大的误差，期望能够每个结果都为30.0
std::cout << "最大误差：" << maxError << std::endl;
}
//释放设备内存
cudaFree(d_x);//释放d_x的内存
cudaFree(d_y);//释放d_y的内存
cudaFree(d_z);//释放d_z的内存
//释放主机内存
free(x);//释放x的内存
free(y);//释放y的内存
free(z);//释放z的内存
return 0;
}这里我们的向量大小为1<<20，而block大小为256，那么grid大小是4096，kernel的线程层级结构如下图所示：
在上面的实现中，我们需要单独在host和device上进行内存分配，并且要进行数据拷贝，这是很容易出错的。好在CUDA 6.0引入统一内存（Unified Memory）来避免这种麻烦，简单来说就是统一内存使用一个托管内存来共同管理host和device中的内存，并且自动在host和device中进行数据传输。CUDA中使用cudaMallocManaged函数分配托管内存：
cudaError_t cudaMallocManaged(void **devPtr, size_t size, unsigned int flag=0);
利用统一内存，可以将上面的程序简化如下：
int main()
{
//使用cudaMallocManaged函数分配托管内存
int N1 = 1 << 20;
int n1Bytes = N1 * sizeof(float);
//申请托管内存
float *x1, *y1, *z1;
cudaMallocManaged((void**)&x1, n1Bytes);
cudaMallocManaged((void**)&y1, n1Bytes);
cudaMallocManaged((void**)&z1, n1Bytes);
//初始化数据
for (int i = 0; i < N1; i++)
{
x1[i] = 10.0;
y1[i] = 20.0;
}
//定义kernel的执行配置
dim3 blockSize1(256);
dim3 gridSize1((N1+blockSize.x-1)/blockSize.x);
//执行kernel
add <<<gridSize1, blockSize1>> > (x1,y1,z1,N1);
//同步device保证结果能正确访问。
cudaDeviceSynchronize();
//检查执行结果
float maxError1 = 0.0;
for (int i = 0; i < N1; i++)
{
maxError1 = fmax(maxError1,fabs(z1[i]-30.0));
std::cout << "最大误差：" << maxError1 << std::endl;
}
//释放内存
cudaFree(x1);
cudaFree(y1);
cudaFree(z1);
}cudaDeviceSynchronize()` 是CUDA编程中用于同步的一个函数，它的作用是确保所有之前发出的、与当前设备（GPU）关联的主机（CPU）线程中的CUDA API调用都已完成执行。在CUDA编程模型中，主机和设备是异步执行的，这意味着当主机代码启动一个核函数后，它不会等待该核函数完成就继续执行后续的主机代码。同样，当主机代码执行一个数据从主机传输到设备的操作时，它也不会等待数据传输完成就继续执行。
由于这种异步性，有时我们需要确保在继续执行主机代码之前，之前的核函数或数据传输操作已经完成。这就是`cudaDeviceSynchronize()`的用处。调用`cudaDeviceSynchronize()`会阻塞调用它的主机线程，直到之前启动的所有核函数和/或数据传输操作都完成。
此外，`cudaDeviceSynchronize()`还可以用于错误检查。如果在之前的核函数执行或数据传输中发生了错误，那么在调用`cudaDeviceSynchronize()`时，这个错误会被捕获并返回给主机。
需要注意的是，虽然`cudaDeviceSynchronize()`很有用，但频繁使用它可能会导致性能下降，因为它会阻塞主机线程。因此，在设计CUDA程序时，应仔细考虑何时需要同步，并尽量减少不必要的同步操作。CUDA还提供了其他同步原语，如流（streams）和事件（events），它们提供了更细粒度的同步控制，但使用起来也更复杂。
矩阵乘法实例
最后我们再实现一个稍微复杂一些的例子，就是2个矩阵的乘法，设输入矩阵为A和B，要得到C=A*B。实现思路是每个线程计算C的一个元素值Cij，对于矩阵运算，应该选用grid和bloc为2-D的。参照CUDA编程入门极简教程_cuda编程举例-优快云博客。
2024.11.30继续学习
cuda实现图像二值化
自己不会写代码，进行了AI。搜出来的代码有点问题，我自己改好了。最后实现效果如下：
二值化之前：
二值化之后：
还是比较艰辛，能够顺利的视线一个功能，咱也是入门了，给自己点个赞。下面附上二值化的代码，并且每个代码是什么意思都添加了注释，注意这些代码都是在***.cu：文件中编写的
我使用的头函数包括：
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include<opencv2/opencv.hpp>
#include <stdio.h>
#include <stdlib.h>
#include<iostream>二值化的核函数代码：
//图像二值化核函数
__global__ void binarizeKernel(unsigned char* inputImage, unsigned char* outputImage, int width, int height, unsigned char threshold) {
int x = blockIdx.x * blockDim.x + threadIdx.x; // 计算当前线程的x坐标
int y = blockIdx.y * blockDim.y + threadIdx.y; // 计算当前线程的y坐标
// 确保当前线程在图像范围内
if (x < width && y < height) {
int idx = y * width + x; // 将二维坐标转换为一维索引
outputImage[idx] = (inputImage[idx] < threshold) ? 0 : 255; // 应用二值化逻辑
}
}二值化函数：
//使用cuda实现图像二值化
void binary(cv::Mat& image,cv::Mat& binary, unsigned char threshold)
{
int width = image.cols;
int height = image.rows;
//将图像转化为一维数组
unsigned char *h_image = image.data;
unsigned char *h_binary = new unsigned char[width*height];
unsigned char *d_image, *d_binary;
int size = width*height * sizeof(unsigned char);
//申请CUDA设备内存
cudaMalloc((void**)&d_image, size);
cudaMalloc((void**)&d_binary, size);
//将输入图像从主机复制到设备
cudaMemcpy(d_image, h_image, size, cudaMemcpyHostToDevice);
//定义内核执行配置
dim3 blockSize(32, 32);
dim3 gridSize((width+blockSize.x-1)/blockSize.x,(height+blockSize.y-1)/blockSize.y);
//执行二值化内核
binarizeKernel <<<gridSize, blockSize >>> (d_image, d_binary, width, height, threshold);
//同步设备
cudaDeviceSynchronize();
//将结果从设备复制回主机
cudaMemcpy(h_binary, d_binary,size,cudaMemcpyDeviceToHost);
//创建图像并保存
cv::Mat outputimage(height,width, CV_8UC1, h_binary);
outputimage.copyTo(binary);
//释放CUDA内存
cudaFree(d_image);
cudaFree(d_binary);
delete[] h_binary;
std::cout << "图像二值化完成，结果已返回。" << std::endl;
}执行主函数：
int main（）
{
cv::Mat image2 = cv::imread("leftlaser1_27.jpg", cv::IMREAD_GRAYSCALE);
cv::Mat binaryimage(image2.rows, image2.cols, CV_8UC1, cv::Scalar(0));
unsigned char a = 180;
binary(image2, binaryimage, a);
return 0;
}
最后在使用image watch看看局部放大的二值化效果，即非0就255，即黑白。
下一步计划学习使用cuda实现灰度重心法提取激光中心线。
扫一扫
2638
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
