GPU高性能编程 CUDA实战 (二) CUDA C 简介

目标

1. 编写第一段CUDA C代码
2. 了解主机(Host)端编写的代码和设备端(Device)编写的代码的区别
3. 如何从主机上运行设备端代码
4. 了解如何在支持CUDA的设备上使用内存
5. 了解如何查询系统中支持CUDA的设备信息

第一个CUDA C程序

CUDA C在很大程度上与标准C没有区别。在GPU上执行的函数通常称为核函数。

#include <iostream>
__global__ void kernel(void) {
}

int main()
{
    kernel<<<1,1>>>();
    printf("Hello,World\n");
    return 0;
}

这个程序中，有一个global修饰的核函数kernel(),调用它时，用了修饰字符<<<1,1>>>，global字符表示这个函数在设备端运行而不是在主机端。
代码来源于书籍

#include <iostream>
__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)));
    add<<<1,1>>>(2,7,dev_c);
    HANDLE_ERROR(cudaMemcpy(&c,dev_c,sizeof(int),cudaMemcpyDeviceTohost));
    printf("2+7=%d\n",c);
    cudaFree(dev_c);
    return 0;
}

在设备端进行运算操作时，有要保存的数据，必须要先在设备端分配空间，然后运算完之后再传输回主机端保存。记住要释放设备端空间。设备端cudaMalloc返回的指针不一定能进行引用。
cudaMalloc()函数是分配设备端内存的，这个函数类似于malloc()，第一个参数是一个指针，指向保存变量的地址，第二个参数是大小。HANDLE_ERROE是自定义的宏，判断是否出错。cudaMemcpy()将设备端数据复制到主机端。
总结如下：
1. 可以将cudaMalloc分配的指针传递给设备上执行的函数。
2. 可以在设备代码中使用cudaMalloc分配的指针进行操作。
3. 可以将cudaMalloc分配的指针传递给主机上执行的函数。
4. 不能在主机代码中使用cudaMalloc分配的指针进行内存读写操作。

访问设备内存两种最常见的方法就是在设备代码中使用设备指针和调用cudaMemcpy函数。

cudaMemcpy函数类似于memcpy函数，只是在多了一个参数来指定设备内存究竟是源指针还是目标指针，它可以是cudaMemcpyDeviceToHost,cudaMemcpyHostToDevice和cudaMemcpyDeviceToDevice。

查询设备

利用CUDA C 可以查询到的信息

设备属性	描述
char name[256]	标识设备的ASCII字符串
size_t totalGlobalMem	设备上全局内存的总量(字节)
size_t sharedMemPerBlock	线程块中可以是哟你哦个的最大共享内存数量(字节)
int regsPerBlock	每个线程中可以用的32位寄存器数量
int warpSize	在一个线程束中包含的线程数量
size_t memPitch	在内存复制中最大的修正量(字节)
int amxThreadsPerBlock	在一个线程块中可以包含最大的线程数量
int maxThreadsDim	多维线程块数组中，每一维可以包含的最大线程数量
size_t totalConstMem	常量内存的总量
int major	设备计算功能集的主版本号
int minor	设备功能集的次版本号
size_t textureAlignment	设备的纹理对齐要求
int deviceOverlap	一个布尔类型的值，表示设备是否可以同时执行一个cudaMemory()调用和一个核函数调用
int multiProcessorCount	设备上多处理器的数量
int kernelExecTimeoutEnabled	一个布尔值，表示在该设备上执行的核函数是否存在运行时限制
int integrated	表示设备是否是一个继承GPU
int canMapHostMemory	一个布尔类型，表示设备是否将主机内存映射到CUDA设备地址空间
int computeMode	表示设备的计算模式：默认（default），独占（Exclusive），禁止（Prohibited）
int maxTexture1D	一维纹理的最大大小
int maxTexture2D[2]	二维纹理的最大维数
int maxTexture3D[3]	三维纹理的最大维数
int maxTexture2DArray[3]	二维纹理数组的最大维数
int concurrentKernels	表示设备是否支持在同一个上下文中同时执行多个核函数

设备查询代码

#include <stdio.h>
int main()
{
    cudaDeviceProp prop;
    int count;
    HANDLE_ERROR(cudaGetDeviceCount(&count));
    for(int i = 0;i<count;i++)
    {
        HANDLE_ERROR(cudaGetDeviceProperties(&prop,i));
        //对设备数量执行某些操作
    }
    return 0;
}

欢迎评论，一起学习。