CUDA常见函数(小白入门)
示例代码0:VS自带Kernel.cu
初始定义
程序首先定义了一个函数addWithCuda,它是调用GPU运算的入口函数,返回类型是cudaError_t。
cudaError_t是一个枚举类型,可以作为几乎所有CUDA函数的返回类型,用来检测函数执行期间发生的不同类型的错误,一共有80多个错误类型,可以在driver_types.h头文件中查看每一个整型对应的错误类型,如果返回0,代表执行成功。
核函数
函数addKernel在最前有一个修饰符“global”,这个修饰符告诉编译器,被修饰的函数应该编译为在GPU而不是在CPU上运行,所以这个函数将被交给编译设备代码的编译器——NVCC编译器来处理,其他普通的函数或语句将交给主机编译器处理。
这里“设备”的概念可以理解为GPU和其显存组成的运算单元,“主机”可以理解为CPU和系统内存组成的运算单元。在GPU上执行的函数称为核函数。
注:关于threadIdx.x的说明。
CUDA中的线程(thread)是设备中并行运算结构中的最小单位,类似于主机中的线程的概念,thread可以以一维、二维、三维的形式组织在一起,threadIdx.x表示的是thread在x方向的索引号,还可能存在thread在y和z方向的索引号threadIdx.y和threadIdx.z。
一维、二维或三维的thread组成一个线程块(Block),一维、二维或三维的线程块(Block)组合成一个线程块网格(Grid),线程块网格(Grid)可以是一维或二维的。通过网格块(Grid)->线程块(Block)->线程(thread)的 顺序可以定位到每一个并且唯一的线程。
addKernel函数会被GPU上的多个线程同时执行一次,线程间彼此没有通信,相互独立。到底会有多少个线程来分别执行核函数,是在“<<< >>>”符号里定义的。“<<< >>>”表示运行时配置符号,在本程序中的定义是<<<1,size>>>,表示分配了一个线程块(Block),每个线程块有分配了size个线程,“<<<>>>”中的 参数并不是传递给设备代码的参数,而是定义主机代码运行时如何启动设备代码。以上定义的这些线程都是一个维度上的,可以通过thredaIdx.x来获取执行当前计算任务的线程的ID号。
其他函数
cudaSetDevice函数用来设置要在哪个GPU上执行,如果只有一个GPU,设置为cudaSetDevice(0);
cudaGetLastError函数用于返回最新的一个运行时调用错误,对于任何CUDA错误,都可以通过函数cudaGetErrorString函数来获取错误的详细信息。
cudaDeviceSynchronize函数提供了一个阻塞,用于等待所有的线程都执行完各自的计算任务,然后继续往下执行。
cudaFree函数用于释放申请的显存空间。
cudaDeviceReset函数用于释放所有申请的显存空间和重置设备状态;
示例代码1:传递参数
#include <iostream>
#include "book.h"
__global__ void add( int a, int b, int *c ) {
*c = a + b;
}
int main( void ) {
int c;
int *dev_c;
HANDLE_ERROR( cudaMalloc( (void**)&dev_c, sizeof(int) ) ); /* 第一个参数是指针,只想用于保存新分配内存地址的变量,第二个参数是分配内存大小 */
/* 返回值不是指针,且返回类型为void*。 */
add<<<1,1>>>( 2, 7, dev_c ); // <<<>>>表示运行时配置符号,第一个参数表示只分配一个线程组(又称线程块、Block),第二个参数表示每个线程组有size个线程(Thread)
HANDLE_ERROE( cudaMemcpy( &c,
dev_c,
sizeof(int)
cudaMemcpyDeviceToHost ) );
print( "2 + 7 = %d\n", c );
cudaFree( dev_c );
return 0;
}
cudaMalloc
memory allocate 动态分配内存
malloc
例:
int *p = (int *)malloc(sizeof(int))
int* p代表一个以int类型地址为内容的指针变量,p这个变量占4个字节(某些计算机),这个p变量是静态分配的一个变量。
malloc函数返回开辟空间的首地址,加(int *)的目的是让计算机知道,如何去划分这个开辟的空间,因为char、int 、long这些类型的字节大小是不一样的,我们知道了首地址,还要知道是以几个字节为单元。
cudaFree
需要使用此函数释放cudaMalloc函数分配的内存。
cudaMemcpy
在主机代码中调用cudaMemcpy()来访问设备上的内存。
cudaMemcpy(目的地, 复制对象, 复制区域长度, 设备到主机/主机到设备)
示例代码1中,cudaMemcpy()中最后一个参数为cudaMemcpyDeviceToHost,意为运行时源指针为设备指针,目标指针为主机指针。
Memcpy
如果源指针和目标指针都在主机上,可直接调用此函数。
例:
void *memcpy(void *dest, const void *src,unsigned int count);
dest - 目的
src - 源
count - 复制字节数
由src所指内存区域复制count个字节到dest所指内存区域。
返回指向dest的指针。
当dest <= src-count 或dest >= src+count时,不会产生覆盖问题,即源数据不会被更改。
若不在以上范围内,则源数据会被更改。
示例代码2:查询设备
#include "../common/book.h"
int main void( void ) {
cudaDeviceProp prop;
int count;
HANDLE_ERROR( cudaGetDevicecount( &count ) );
for (int i=0; i< count; i++){
HANDLE_ERROR( cudaGetDeviceProperties( &prop, i ) );
// 对设备属性执行某些操作
}
}
int count;
HANDLE_ERROR( cudaGetDeviceCount( &count ) );/* 获取CUDA设备的数量。 */
cudaGetDeviceCount
在调用此函数后,可以对每个设备进行迭代,并查询各个设备相关信息。
CUDA运行时将返回一个cudaDeviceProp类型的结构,其中包含了设备的相关属性。
参考《NVIDIA CUDA Programming Guide》
#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
using namespace std;
int main()
{
cudaDeviceProp deviceProp;
int deviceCount;
cudaError_t cudaError;
cudaError = cudaGetDeviceCount(&deviceCount);
for (int i = 0; i < deviceCount; i++)
{
cudaError = cudaGetDeviceProperties(&deviceProp, i);
cout << "设备 " << i + 1 << " 的主要属性: " << endl;
cout << "设备显卡型号: " << deviceProp.name << endl;
cout << "设备上一个线程束(Warp)中包含的线程数量: " << deviceProp.warpSize << endl;
cout << "多维线程块(Block)数组中,每一维可包含的最大线程(Thread)数量: " << deviceProp.maxThreadsDim[3] << endl;
cout << "一个线程格(Grid)中,每一维可以包含线程块(Block)数量: " << deviceProp.maxGridSize[3] << endl;
cout << "设备全局内存总量(以MB为单位): " << deviceProp.totalGlobalMem / 1024 / 1024 << endl;
cout << "设备上一个线程块(Block)中可用的最大共享内存(以KB为单位): " << deviceProp.sharedMemPerBlock / 1024 << endl;
cout << "设备上一个线程块(Block)种可用的32位寄存器数量: " << deviceProp.regsPerBlock << endl;
cout << "设备上一个线程块(Block)可包含的最大线程数量: " << deviceProp.maxThreadsPerBlock << endl;
cout << "设备的计算功能集(Compute Capability)的版本号: " << deviceProp.major << "." << deviceProp.minor << endl;
cout << "设备上多处理器的数量: " << deviceProp.multiProcessorCount << endl;
}
getchar();
return 0;
}
其他
快速傅里叶变换(C++)
floor
floor函数,其功能是“向下取整”,或者说“向下舍入”、“向零取舍”,即取不大于x的最大整数
ceil
double ceil(double x);
float ceil ( float value );
功 能: 返回大于或者等于指定表达式的最小整数,value 如果有小数部分则进一位。
头文件math.h
gettimeofday(windows下的实现)
#include <time.h>
#ifdef WIN32
#include <windows.h>
#else
#include <sys/time.h>
#endif
#ifdef WIN32
int gettimeofday(struct timeval *tp, void *tzp)
{
time_t clock;
struct tm tm;
SYSTEMTIME wtm;
GetLocalTime(&wtm);
tm.tm_year = wtm.wYear - 1900;
tm.tm_mon = wtm.wMonth - 1;
tm.tm_mday = wtm.wDay;
tm.tm_hour = wtm.wHour;
tm.tm_min = wtm.wMinute;
tm.tm_sec = wtm.wSecond;
tm.tm_isdst = -1;
clock = mktime(&tm);
tp->tv_sec = clock;
tp->tv_usec = wtm.wMilliseconds * 1000;
return (0);
}
#endif