最近3个星期都没有更新文章了，花全部心思为了完成上面布置的任务了，需求是对指纹数据进行插值和截取操作，实现算法优化，之前的代码是直接暴力循环的，对于扫描指纹和出图整个过程花费时间25秒左右，和目前最快的速度13秒左右相差太大，对于指纹识别商用肯定不切实际的，因此打算开始优化，首先开始测试每个步骤所花费的时间，经过师兄的测试，发现插值和截取所花费的时间占大头，因此导师把这个任务交给了我，我瑟瑟发抖，毕竟对底层c++加速了解好少，谁叫是研究僧呢，肯定啥都要研究，话不多说，开干，我开始首先配置好cuda环境，这个好弄，我之前用python的tensorflow框架就搞好了，然后就是vscode配置cuda文件代码运行环境了，因为.cu后缀文件支持c++代码的编写了，可以直接调用cuda命令运行，那就好办了，先了解下cuda中gpu并行的知识，一般我们代码没特别函数或者声明都是默认再cpu运行的，在cuda中，有个__global__函数，也叫核函数，是在gpu上线程中并行执行的函数，注意的是这个函数没有返回值，一般类型都设置为void，先看下以下代码
#include<stdio.h>
#include<stdlib.h>
#include<cuda_runtime_api.h>
#include<iostream>
__global__ void kennel(int arr[3][3],int n)
{
int i = threadIdx.x+blockDim.x*blockIdx.x;
int j = threadIdx.y+blockDim.y*blockIdx.y;
for(i;i<3;i+=gridDim.x * blockDim.x)
{
for(j;j<3;j+=gridDim.y * blockDim.y)
{
printf("i:->%d,j:->%d...%d\n",i,j,arr[i][j]);
}
}
}
int main()
{
int N = 9;
float aaa = 3/10;
int datas[3][3] = {{1,2,3},{4,5,6},{7,8,9}};
size_t size = N * sizeof(int);
int (*arr)[3];//定义指针指向n行3列的数组
cudaMallocManaged(&arr,N*sizeof(int));
cudaMemcpy(arr,datas,size,cudaMemcpyHostToDevice);
kennel<<<3,3>>>(arr,N);
cudaDeviceSynchronize();
return 0;
}
看上去很复杂，其实就是对3行3列的二维数组进行3*3线程并行输出，使用命令运行下，如图所示
PS C:\Users\14499\Desktop\vscodeProject> nvcc .\demo3.cu -o .\demo3
demo3.cu
./demo3.cu : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(825) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(1760) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(2622) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(3461) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(4404) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(5302) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(6213) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(7094) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(7893) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt/device_functions.h : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt/device_functions.h(773) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt/device_functions.h(1629) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\device_double_functions.h : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\sm_20_intrinsics.h : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\sm_20_intrinsics.h(926) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
./demo3.cu(10): warning: expression has no effect
./demo3.cu(12): warning: expression has no effect
./demo3.cu : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(825) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(1760) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(2622) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(3461) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(4404) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(5302) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(6213) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(7094) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(7893) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt/device_functions.h : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt/device_functions.h(773) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt/device_functions.h(1629) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\device_double_functions.h : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\sm_20_intrinsics.h : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\sm_20_intrinsics.h(926) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
./demo3.cu(10): warning: expression has no effect
./demo3.cu(12): warning: expression has no effect
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(825) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(1760) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(2622) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(3461) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(4404) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(5302) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(6213) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(7094) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
c:\program files\nvidia gpu computing toolkit\cuda\v10.0\include\crt\math_functions.h(7893) : warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
正在创建库 ./demo3.lib 和对象 ./demo3.exp
PS C:\Users\14499\Desktop\vscodeProject> .\demo3
i:->0,j:->0...1
i:->1,j:->0...4
i:->2,j:->0...7
i:->0,j:->1...2
i:->1,j:->1...5
i:->2,j:->1...8
i:->0,j:->2...3
i:->1,j:->2...6
i:->2,j:->2...9
警告可以忽略，这里注意调用核函数要用<<<线程块数,每个线程块对应的线程数>>>(参数一，参数二)表示，除此之外要严格区分cpu上的变量和gpu中的变量，gpu和cpu中的变量不可直接互相调用，对于gpu，必须动态分配gpu内存来存储变量，然后将cpu中的变量拷贝到gpu中，这样就可以实现cpu和gpu之间变量的传递了。还有线程块*线程数就是总线程数了，这里代码中的gridDim.x为3，blockDim.x也为3.总线程就是9个。了解了这个知识后，就开始实现截取和插值的算法优化了，直接上代码，
#include<stdio.h>
#include<fstream>
#include<cuda_runtime_api.h>
#include<iostream>
#include <string>
#include <cstdlib>
#include<cstring>
using namespace std;
int datas[1024000];
int newData[500*1500];
float binData[4096];
//插值
float allendDatas[500*4096];
//读取bin文件
void readBin()
{
//打开标定文件
std::ifstream inF("3001800cali.bin", std::ios::binary);
inF.read((char*)binData, sizeof(float) * (4096));
inF.close();
for(int i=0;i<4096;i++)
{
printf("%f\n",binData[i]);
}
}
//读取txt数据
void readData(int N)
{
ifstream infile("100.txt",ios::in);
int i=0;
char data[10];
while(infile.getline(data, 10))
{
int a = atoi(data);
datas[i] = a;
printf("this is %d nums:%d\n",i,a);
i++;
}
}
// CUDA 截取核函数，
__global__ void Speedjiequ(int *p,int *news,int start,int end) {
int i;
int j;
int new_index;
int index;
int idx;
int idy;
for(idy=blockIdx.y * blockDim.y + threadIdx.y;idy<500;idy += blockDim.y * gridDim.y)
{
for(idx=blockIdx.x *blockDim.x + threadIdx.x;idx<2048;idx+= blockDim.x * gridDim.x)
{
if(idx>=start && idx<end)
{
index = idy*2048+idx;
new_index = index+idy*(end-start)-idy*2048-start;
news[new_index] = p[index-1];
// printf("new_index:%d\n",new_index);
// printf("********idy:%d,idx:%d***value is:%d***index:%d\n",idy,idx,p[index-1],index-1);
}
}
}
}
//cuda 插值核函数
__global__ void Speedchazhi(int * pp,float * endDatas,float *BinData)
{
int idx;
int idy;
float k;
float b;
for(idy=blockIdx.y * blockDim.y + threadIdx.y;idy<500;idy += blockDim.y * gridDim.y)
{
for(idx=blockIdx.x *blockDim.x + threadIdx.x;idx<4096;idx+= blockDim.x * gridDim.x)
{
//binData为一维数组；
float bin_index = BinData[idx];
if ((int)bin_index < 1499)
{
k = pp[(int)(1500*idy+bin_index+1)]-pp[(int)(1500*idy+bin_index)];
b = pp[(int)(bin_index+1500*idy)]- k*(int)(bin_index+1500*idy);
// printf("k is %f,b is %f\n",k,b);
endDatas[idy*4096+idx] = k*(bin_index+1500*idy)+b;
}
else
{
k = pp[(int)(bin_index+1500*idy)]-pp[(int)(bin_index-1+1500*idy)];
b = pp[(int)(bin_index-1+1500*idy)]-k*(int)(bin_index-1+1500*idy);
endDatas[idy*4096+idx] = k*(bin_index+1500*idy)+b;
}
}
}
}
int main() {
//设置显卡设备
readBin();
cudaSetDevice(0);
int N = 1024000;
int n = 500*1500;
int *gpuDatas;
int *newGpuDatas;
//插值数据
float *endDatas;
int *newImgs;
float *gpuBinData;
// int *newImgs;
time_t start,stop,start1,stop1,start2,stop2;
size_t size = N * sizeof(int);
size_t size1 = n * sizeof(int);
size_t size2 = 2*N * sizeof(float);
size_t size3 = 4096* sizeof(float);
//截取数据
cudaMallocManaged(&gpuDatas, size); // 为a分配CPU和GPU空间
cudaMallocManaged(&newGpuDatas, size1); // 为a分配CPU和GPU空间
//插值数据
cudaMallocManaged(&endDatas, size2);
cudaMallocManaged(&newImgs, size1);
cudaMallocManaged(&gpuBinData, size3);
// 将数据从内存拷贝到显存
start = time(NULL);
readData(N); // 为数组a赋值
stop = time(NULL);
printf("read txt_time is:%d\n",stop-start);
//截取数据拷贝到显存
cudaMemcpy(gpuDatas,datas,size,cudaMemcpyHostToDevice);
cudaMemcpy(newGpuDatas,newData,size1,cudaMemcpyHostToDevice);
//插值数据拷贝到显存
cudaMemcpy(endDatas,allendDatas,size2,cudaMemcpyHostToDevice);
cudaMemcpy(gpuBinData,binData,size3,cudaMemcpyHostToDevice);
// cudaMemcpy(newImgs,imgs,size1,cudaMemcpyHostToDevice);
start1 = time(NULL);
//截取函数
Speedjiequ<<<1000,1024>>>(gpuDatas,newGpuDatas,300,1800); // 执行核函数
stop1 = time(NULL);
printf("gpu_time_is:%d\n",stop1-start1);
// 将截取数据从显存拷贝到内存并循环输出
cudaMemcpy(newData,newGpuDatas,n*sizeof(int),cudaMemcpyDeviceToHost);
// ofstream out("jiequ.txt");
// for(int i=0;i<500;i++)
// {
// for(int j=0;j<1500;j++)
// {
// out << newData[i*1500+j] << endl;
// printf("alfter jiequ data is:---%d-----\n",newData[i*1500+j]);
// }
// }
//插值核函数运行
cudaMemcpy(newImgs,newData,size1,cudaMemcpyHostToDevice);
start2 = time(NULL);
Speedchazhi<<<2000,1024>>>(newImgs,endDatas,gpuBinData);
stop2 = time(NULL);
printf("chazhi_time_is:%d\n",stop1-start1);
//将插值数据从显存拷贝到内存并打印
cudaMemcpy(allendDatas,endDatas,size2,cudaMemcpyDeviceToHost);
ofstream out("chazhi.txt");
for(int i=0;i<500;i++)
{
for(int j=0;j<4096;j++)
{
out << allendDatas[i*4096+j] << endl;
printf("alfter chazhi data is:xxxxxxx%fxxxxxxx\n",allendDatas[i*4096+j]);
}
}
out.close();
cudaDeviceSynchronize(); // 同步
//释放所有显存开辟的空间
cudaFree(gpuDatas);
cudaFree(newGpuDatas);
cudaFree(endDatas);
cudaFree(gpuBinData);
cudaDeviceReset();
printf("gpu_time_is:%d\n",stop1-start1);
return 0;
}
3个星期的结果，我也写吐了，差点放弃了，记录实现后发现对百万数据进行操作，时间基本为0，效率直接提升，原来之前的所用时间为12秒，现在直接节省12秒，只能说gpu并行牛皮，c++牛皮，多谢导师鼓励和给我锻炼的机会，多谢师兄的耐心详细的流程讲解，本次仅为记录，我是菜鸡一枚。
如果上述代码帮助您很多，可以打赏下以减少服务器的开支吗，万分感谢！
赣公网安备 36092402000079号
点击此处登录后即可评论