GPU高性能编程CUDA实战-第三章

CPU以及系统内存称为主机，GPU及其内存称为设备。

在GPU上执行的函数称为核函数(Kernel)。

C的malloc(),memcpy()和free()等API对应CUDA API：cudaMalloc()分配设备内存,cudaMemcpy()在设备和主机之间复制数据，cudaFree()释放设备内存。

//核函数调用
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <iostream>

__global__ void kernel(void) {      //__global__修饰符告诉编译器，函数应该编译为在设备而不是主机运行
}                                   //函数kernel()将被交给编译设备代码的编译器

int main(void) {                    //main()函数将被交给主机编译器
    kernel <<<1, 1 >>>();           //在主机代码中调用设备代码，尖括号表示要将一些参数传递给运行时系统，告诉运行时如何启动设备代码
    printf("Hello World!\n");
    return 0;
}

book.h：(15条消息) 《GPU高性能编程》——book.h_bmt兔子的博客-CSDN博客

//传递参数
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>

__global__ void add(int a, int b, int *c)
{
    *c = a + b;    //将参数a和b相加，并将结果保存在c指向的内存中
}

int main(void)
{
    int c;
    int* dev_c;
    HANDLE_ERROR(cudaMalloc((void**)&dev_c, sizeof(int)));    //通过cudaMalloc()函数告诉CUDA运行时在设备上分配内存
                                                              //第一个参数是一个指针，指向用于保存新分配内存地址的变量；第二个参数是分配内存的大小
    add << <1, 1 >> > (2, 7, dev_c);
//程序员一定不能在主机代码中对cudaMalloc()返回的指针进行解引用
//主机代码可以通过调用cudaMemcoy()来访问设备上的内存
    HANDLE_ERROR(cudaMemcpy(&c,   //函数调用外层的HANDLE_ERROR()是定义的一个宏，这个宏只是判断函数调用是否返回了一个错误值
                            dev_c,    
                            sizeof(int),
                            cudaMemcpyDeviceToHost));    //cudaMemcpyDeviceToHost表示运行时源指针是一个设备指针，目标指针是一个主机指针
    printf("2+7=%d\n", c);                               //cudaMemcpyHostToDevice表示运行时源指针是一个主机指针，目标指针是一个设备指针
    cudaFree(dev_c);                                     //cudaMemcpyDeviceToDevice表示运行时两个指针都位于设备上；如果源指针和目标指针都位于主机上，可以直接调用C的memcpy()函数

    return 0;
}
//设备指针的使用限制：
//可以将cudaMalloc()分配的指针传递给在设备上执行的函数
//可以在设备代码中使用cudaMalloc()分配的指针进行内存读写操作
//可以将cudaMalloc()分配的指针传递给在主机上执行的函数
//不能再主机代码中使用cudaMalloc()分配的指针进行内存读写操作
//推论：不能使用标准C的free()函数来释放cudaMalloc()分配的内存，要释放cudaMalloc()分配的内存，需要调用cudaFree()
//主机指针只能访问主机代码中的内存，设备指针也只能访问设备代码中的内存

//查询设备
int count;
HANDLE_ERROR(cudaGetDeviceCount(&count));    //获得CUDA设备的数量

struct cudaDeviceProp{    //获得设备相关属性
    char name[256];
    size_t totalGlobalMem;
    ......
}

//一些具体操作
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>

int main(void) {
    cudaDeviceProp prop;
    
    int count;
    HANDLE_ERROR(cudaGetDeviceCount(&count));
    for (int i = 0; i < count; i++) {
        HANDLE_ERROR(cudaGetDeviceProperties(&prop,i));
        printf("---General Information for device %d ---\n", i);
        printf("Name: %s\n", prop.name);
        printf("Compute capability: %d.%d\n", prop.major, prop.minor);
        printf("Clock rate: %d\n", prop.clockRate);
        printf("Device copy overlap: ");
        if (prop.deviceOverlap)
            printf("Enabled\n");
        else
            printf("Disabled\n");
        printf("Kernel execition timeout:");
        if (prop.kernelExecTimeoutEnabled)
            printf("Enabled\n");
        else
            printf("Disabled\n");

        printf("---Memory Information for device %d ---\n", i);
        printf("Total global mem: %ld\n", prop.totalGlobalMem);
        printf("Total constant Mem: %ld\n", prop.totalConstMem);
        printf("Max mem pitch: %ld\n", prop.memPitch);
        printf("Texture Alignment: %ld\n", prop.textureAlignment);
        printf("---MP Information for device %d---\n", i);
        printf("Multiprocessor count: %d\n", prop.multiProcessorCount);
        printf("Shared mem per mp: %ld\n", prop.sharedMemPerBlock);
        printf("Registers per mp: %d\n", prop.regsPerBlock);
        printf("Threads in warp: %d\n", prop.warpSize);
        printf("Max threads per block: %d\n", prop.maxThreadsPerBlock);
        printf("Max thread dimensions: (%d,%d,%d)\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
        printf("\n");
    }
}

//设备属性的使用：应用程序从多个GPU中选择最合适的GPU
//找出希望设备拥有的属性并将这些属性填充到一个cudaDeviceProp结构
cudaDeviceProp prop;
memset(&prop, 0, sizeof(cudaDeviceProp));
prop.major = 1;
prop.minor = 3;


int main(void) {
    cudaDeviceProp prop;
    int dev;

    HANDLE_ERROR(cudaGetDevice(&dev));    
    printf("ID of current CUDA device: %d\n", dev);

    memset(&prop, 0, sizeof(cudaDeviceProp));
    prop.major = 1;
    prop.minor = 3;
    Handle_ERROR(cudaChooseDevice(&dev, &prop));    //填充完cudaDeviceProp结构后，将其传递给cudaChoiceDevice(),CUDA运行时将查找是否存在某个设备满足这些条件
    printf("ID of CUDA device closest to revision 1.3: %d\n", dev);    //cudaChooseDevice()函数将返回一个设备ID，可以将这个ID传递给cudaSetDevice(),随后所有的设备操作都将在这个设备上执行
    HANDLE_ERROR(cudaSetDevice(dev));
}