CUDA 入门笔记-1 -向量相加，计时

最新推荐文章于 2023-09-20 14:46:39 发布

SIXtong

最新推荐文章于 2023-09-20 14:46:39 发布

阅读量440

点赞数

分类专栏： cuda 文章标签： Cuda linux

本文链接：https://blog.csdn.net/SIXtong/article/details/102556345

版权

cuda 专栏收录该内容

4 篇文章 0 订阅

订阅专栏

来源CUDA C权威编程指南

1 内存管理

内存分为两种：
CPU及其内存-主机内存
GPU及其内存-设备内存

cuda c的内存管理类似于C语言。

功能	C	CUDA
申请内存	malloc	cudaMalloc
初始化内存	memset	cudaMemset
申请内存	free	cudaFree
复制	memcpy	cudaMemcpy

cudaError_t cudaMalloc (void **devPtr, size_t size );
对于cudaMalloc，注意参数的格式，是双重指针。

如果已分配的主机内存A，已分配的主机内存B之间直接赋值A=B会发生错误(cuda 6.0之前)。
需要使用cudaMemcpy(A，B，nBytes，cudaMemcpyDeviceToHost)；
其中nBytes是空间的大小，单位为字节，最后一个参数为类型kind，有四种。
1 cudaMemcpyDeviceToHost - GPU->CPU
2 cudaMemcpyHostToDevice - CPU->GPU
3 cudaMemcpyDeviceToDevice - GPU->GPU
4 cudaMemcpyHostToHost - CPU->CPU

cuda 6.0之后提出统一寻址。

cuda采用线程块和线程网络来管理线程。坐标变量有如下两种：
·blockIdx(线程块在线程格内的索引)
·threadIdx(块内的线程索引)
在这里插入图片描述由一个内核启动所产生的所有线程统称为一个网格。同一网格中的所有线程共享相同的全局内存空间。一个网格由多个线程块构成,一个线程块包含一组线程,同一线程块内的线程协作可以通过·同步·共享内存来实现。
不同块内的线程不能协作。

通常,一个线程格会被组织成线程块的二维数组形式,一个线程块会被组织成线程的三维数组形式。

示例代码

#include <stdio.h>

__global__ void checkIndex(){
	printf("threadIdx(%d,%d,%d) blockIdx(%d,%d,%d) blockDim(%d,%d,%d) gridDim(%d,%d,%d)\n",
			threadIdx.x,threadIdx.y,threadIdx.z,blockIdx.x,blockIdx.y,blockIdx.z,
			blockDim.x,blockDim.y,blockDim.z,gridDim.x,gridDim.y,gridDim.z);
}
//打印坐标和size

int main(){
	int nElem = 6;//线程总数
	dim3 block(3);//dim3类型，3*1*1，用作线程块
	dim3 grid((nElem+block.x-1)/block.x);//dim3类型，2*1*1，用作线程网格，包含两个block
	//网格大小是块大小的倍数，块大小×网格大小 = 总线程数

	printf("grid.x:%d grid.y:%d grid.z:%d\n",grid.x,grid.y,grid.z);
	printf("block.x:%d block.y:%d block.z:%d\n",block.x,block.y,block.z);

	checkIndex<<<grid,block>>>();
	//grid，有几个启动块，每个启动块有几个线程。
	cudaDeviceReset();//清除空间

	return 0;

cuda的同步异步

核函数的调用与主机线程是异步的，也就是核函数调用完成后，控制权马上返还给主机端CPU。但是可以通过下列函数强制主机端等待所有核函数执行完毕
cudaError_t cudaDeviceSynchronize(void);
一些cuda运行时API在主机和设备之间是隐式同步的。如当使用cudaMemcpy函数进行拷贝的时候，主机端必须等数据拷贝完成才能继续执行。

编写核函数

核函数是在设备端执行的代码。
在这里插入图片描述核函数有以下局限：
1、只能访问设备内存
2、必须有void返回类型
3、不支持可变数量的参数
4、不支持静态变量
5、显示异步行为

CUDA的时间

下面的函数返回从1970年1月1日起，到现在经过的秒数，即系统时间。

double cpuSecond(){
	struct timeval tp;
	gettimetoday(&tp,NULL);
	//返回 秒+微秒
	return (double)tp.tv_sec +(double )tp.tv_usec*1e-6;
}

可以使用cuda自带的工具nvprof来进行计时，使用方式如下。
在这里插入图片描述

下面的程序使用nvprof工具

$ nvprof ./main

得到的结果为：

./main Starting ......
==8383== NVPROF is profiling process 8383, command: ./main
Using Device 0: GeForce 940MX
Vector size:16777216
Init time elapsed:0.744784sec
sumArraysOnGPU<<<65536,256>>> Time elapsed 0.013847 sec
Host time elapsed:0.082980sec
Arrays match!
==8383== Profiling application: ./main
==8383== Profiling result:
            Type  Time(%)      Time     Calls       Avg       Min       Max  Name
 GPU activities:   51.47%  57.525ms         2  28.762ms  28.434ms  29.091ms  [CUDA memcpy HtoD]
                   36.65%  40.958ms         1  40.958ms  40.958ms  40.958ms  [CUDA memcpy DtoH]
                   11.88%  13.271ms         1  13.271ms  13.271ms  13.271ms  sumArraysOnGPU(float*, float*, float*, int)
      API calls:   49.53%  122.75ms         3  40.917ms  196.74us  122.31ms  cudaMalloc
                   40.03%  99.195ms         3  33.065ms  28.627ms  41.823ms  cudaMemcpy
                    5.57%  13.796ms         1  13.796ms  13.796ms  13.796ms  cudaDeviceSynchronize
                    4.09%  10.125ms         3  3.3749ms  262.63us  4.9409ms  cudaFree
                    0.33%  811.41us        96  8.4520us     245ns  363.41us  cuDeviceGetAttribute
                    0.33%  808.27us         1  808.27us  808.27us  808.27us  cudaGetDeviceProperties
                    0.07%  164.98us         1  164.98us  164.98us  164.98us  cuDeviceTotalMem
                    0.05%  114.83us         1  114.83us  114.83us  114.83us  cuDeviceGetName
                    0.01%  36.699us         1  36.699us  36.699us  36.699us  cudaLaunchKernel
                    0.00%  6.7910us         1  6.7910us  6.7910us  6.7910us  cudaSetDevice
                    0.00%  6.4320us         1  6.4320us  6.4320us  6.4320us  cuDeviceGetPCIBusId
                    0.00%  2.6050us         3     868ns     246ns  1.9050us  cuDeviceGetCount
                    0.00%  2.4080us         2  1.2040us     307ns  2.1010us  cuDeviceGet
                    0.00%     418ns         1     418ns     418ns     418ns  cuDeviceGetUuid

编写给予GPU的加法的向量函数：
向量相加，计时，比较时间。
块大小256.

#include <stdio.h>
#include <cuda_runtime.h>
#include <sys/time.h>

#define CHECK(call) {\
const cudaError_t error = call;\
	if(error != cudaSuccess){\
		printf("Error: %s:%d, ",__FILE__,__LINE__);\
		printf("code:%d,reason:%s\n",error,cudaGetErrorString(error));\
	}\
}\

double cpuSecond(){
	struct timeval tp;
	gettimeofday(&tp,NULL);
	//秒+微秒
	return (double)tp.tv_sec +(double )tp.tv_usec*1e-6;
}

void checkResult(float *hostRef,float *gpuRef,const int N){
	double epslion = 1e-8;
	bool match = 1;
	for(int i=0;i<N;i++){
		if(abs(hostRef[i]-gpuRef[i])>epslion){
			match = 0;
			printf("Arrays do not match!\n");
			printf("host %5.2f gpu %5.2f at current %d\n",hostRef[i],gpuRef[i],i);
			break;
		}
	}
	if(match) printf("Arrays match!\n");
}

void initData(float *ip,int size){
	time_t t;
	srand((unsigned)time(&t));
	for(int i=0;i<size;i++){
		ip[i] = (float)(rand()&0xFF)/10.0f;
	}
}

void sumArraysOnHost(float *A,float *B,float *C,const int N){
	for(int idx = 0;idx<N;idx++){
		C[idx] = A[idx]+B[idx];
	}
}

__global__ void  sumArraysOnGPU(float *A,float *B,float *C,const int N){
	int i = blockIdx.x*blockDim.x+threadIdx.x;
	//printf("fffffffffffff %d %d %d %d \n",i,C[i],A[i],B[i]);
	C[i] = A[i] + B[i];
}

int main(int argc,char ** argv){
	printf("%s Starting ......\n",argv[0]);
	int dev = 0;
	cudaDeviceProp deviceProp;
	CHECK(cudaGetDeviceProperties(&deviceProp,dev));
	printf("Using Device %d: %s\n",dev,deviceProp.name);
	CHECK(cudaSetDevice(dev));

	int nElem = 1<<24;
	printf("Vector size:%d\n",nElem);
	size_t nBytes = nElem*sizeof(float);
	float *h_A,*h_B,*hostRef,*gpuRef;
	h_A = (float *)malloc(nBytes);
	h_B = (float *)malloc(nBytes);
	hostRef = (float *)malloc(nBytes);
	gpuRef = (float *)malloc(nBytes);

	double iStart,iElaps;

	iStart = cpuSecond();
	initData(h_A,nElem);
	initData(h_B,nElem);
	iElaps = cpuSecond()-iStart;
	printf("Init time elapsed:%lfsec\n",iElaps);

	float *d_A,*d_B,*d_C;
	cudaMalloc((float **)&d_A,nBytes);
	cudaMalloc((float **)&d_B,nBytes);
	cudaMalloc((float **)&d_C,nBytes);

	cudaMemcpy(d_A,h_A,nBytes,cudaMemcpyHostToDevice);
	cudaMemcpy(d_B,h_B,nBytes,cudaMemcpyHostToDevice);

	int iLen = 256;
	dim3 block(iLen);
	dim3 grid((nElem+block.x-1)/block.x);

	iStart = cpuSecond();
	sumArraysOnGPU<<<grid,block>>>(d_A,d_B,d_C,nElem);
	cudaDeviceSynchronize();
	iElaps = cpuSecond()-iStart;
	printf("sumArraysOnGPU<<<%d,%d>>> Time elapsed %lf sec\n",grid.x,block.x,iElaps);

	cudaMemcpy(gpuRef,d_C,nBytes,cudaMemcpyDeviceToHost);

	iStart = cpuSecond();
	sumArraysOnHost(h_A,h_B,hostRef,nElem);
	iElaps = cpuSecond()-iStart;
    printf("Host time elapsed:%lfsec\n",iElaps);

	checkResult(hostRef,gpuRef,nElem);

	cudaFree(d_A);
	cudaFree(d_B);
	cudaFree(d_C);
	free(h_A);
	free(h_B);
	free(hostRef);
	free(gpuRef);

	return 0;
}