本節(jié)目標：

• 為主機（Host）編寫代碼和為設備（Device）編寫代碼的區(qū)別

• 從主機運行設備代碼

• 如何在支持CUDA的設備上使用設備內存

• 如何查詢系統(tǒng)中支持CUDA的設備信息

1. 在GPU設備上執(zhí)行的函數(shù)稱為核函數(shù)，NVIDIA工具使用CUDA C編寫核函數(shù)，CUDA C為標準C增加 global 修飾符，這個修飾符將告訴編譯器，函數(shù)應該編譯為設備而不是主機運行。具有該修飾符的函數(shù)將作為核函數(shù)交給CUDA編譯器編譯。

2.從主機運行設備代碼，如下所示：

cudaError_t err = cudaSuccess;
    ////參數(shù)傳遞
    int c;
    int *dev_c;
    err = cudaMalloc((void**)&dev_c, sizeof(int));
    add_kernel << <1, 1 >> > (2, 7, dev_c); ////核函數(shù)處理
    if (err != cudaSuccess)
    {
        fprintf(stderr, "Failed to allocate device value dev_c (error code %s)!\n", cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }

    err = cudaMemcpy(&c, dev_c, sizeof(int), cudaMemcpyDeviceToHost); ////主機代碼中，不能對cudaMalloc()返回的指針進行讀取，寫入內存操作（即不可修改），但可以作參數(shù)傳遞，執(zhí)行算術運算。
    if (err != cudaSuccess)
    {
        fprintf(stderr, "Failed to copy device value dev_c to c (error code %s)!\n", cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }

    printf("2+7=%d\n", c);
    cudaFree(dev_c); ////不可使用free()，因為在主機代碼中不可修改該內存的值  ，  總的來說，主機指針只能訪問主機代碼的內存，設備指針只能訪問設備代碼的內存

從主機運行設備代碼，以下幾點值得注意：

1.可以像調用C函數(shù)那樣將參數(shù)傳遞給核函數(shù)

2.當設備執(zhí)行任何有用操作時，都需要分配內存，例如將計算值返回主機。

3.使用設備內存：關于cudaMalloc()分配到的指針dev_c.

可以將dev_c傳遞給設備上執(zhí)行的函數(shù)?？梢栽谠O備代碼中使用dev_c進行內存讀/寫操作（修改該內存的值）?？梢钥梢詫?strong>dev_c傳遞給主機上執(zhí)行的函數(shù)，但是不能在主機代碼中對dev_c進行內存讀/寫操作。因為dev_c申請的是設備內存，不能在主機代碼中用主機指針訪問設備內存，反過來也一樣，不能用設備指針訪問主機內存。

然后可以利用cudaMencpy()在主機代碼中訪問設備內存，設備指針用完后用cudaFree()釋放，不可以直接Free(),原因同上，不能在主機代碼中直接訪問設備內存。

4.查詢系統(tǒng)中支持CUDA的設備信息

////查詢設備
    cudaDeviceProp prop;
    int count;
    err = cudaGetDeviceCount(&count);
    for (int i = 0; i < count; i++)
    {
        err = cudaGetDeviceProperties(&prop, i);

        if (err != cudaSuccess)
        {
            fprintf(stderr, "Failed to get device %d properties (error code %s)!\n", i, cudaGetErrorString(err));
            exit(EXIT_FAILURE);
        }

        printf("Name: %s\n", prop.name);////這里打印設備信息，具體包含哪些信息查看《NVIDIA CUDA Programming Guide》  
        printf("Compute capability: %d . %d\n", prop.major,prop.minor);
        printf("Total global Mem:%lld\n", prop.totalGlobalMem);
        printf("Total constant Mem:%ld\n", prop.totalConstMem);

    }

完整代碼：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <helper_cuda.h>
using namespace std;




__global__ void add_kernel(int a,int b, int *c) {
    *c = a + b;
}

int main(void)
{
    cudaError_t err = cudaSuccess;
    ////參數(shù)傳遞
    int c;
    int *dev_c;
    err = cudaMalloc((void**)&dev_c, sizeof(int));
    add_kernel << <1, 1 >> > (2, 7, dev_c); ////核函數(shù)處理
    if (err != cudaSuccess)
    {
        fprintf(stderr, "Failed to allocate device value dev_c (error code %s)!\n", cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }

    err = cudaMemcpy(&c, dev_c, sizeof(int), cudaMemcpyDeviceToHost); ////主機代碼中，不能對cudaMalloc()返回的指針進行讀取，寫入內存操作（即不可修改），但可以作參數(shù)傳遞，執(zhí)行算術運算。
    if (err != cudaSuccess)
    {
        fprintf(stderr, "Failed to copy device value dev_c to c (error code %s)!\n", cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }

    printf("2+7=%d\n", c);
    cudaFree(dev_c); ////不可使用free()，因為在主機代碼中不可修改該內存的值  ，  總的來說，主機指針只能訪問主機代碼的內存，設備指針只能訪問設備代碼的內存

    ////查詢設備
    cudaDeviceProp prop;
    int count;
    err = cudaGetDeviceCount(&count);
    for (int i = 0; i < count; i++)
    {
        err = cudaGetDeviceProperties(&prop, i);

        if (err != cudaSuccess)
        {
            fprintf(stderr, "Failed to get device %d properties (error code %s)!\n", i, cudaGetErrorString(err));
            exit(EXIT_FAILURE);
        }

        printf("Name: %s\n", prop.name);////這里打印設備信息，具體包含哪些信息查看《NVIDIA CUDA Programming Guide》  
        printf("Compute capability: %d . %d\n", prop.major,prop.minor);
        printf("Total global Mem:%lld\n", prop.totalGlobalMem);
        printf("Total constant Mem:%ld\n", prop.totalConstMem);

    }
    return 0;   
    
}