CUDA學習筆記1：第一個CUDA實例

一、cuda簡介

CUDA是支持c++/c語言，一般我喜歡用c來寫，他的編譯是gpu部分由nvcc來進行的

 

一般的函數定義 void  function();

cuda的函數定義 __global__ void function();

 

解釋：在這里，這個global前綴表明這個函數在哪里執行，可以由誰來呼叫

global:主機呼叫，設備執行

host:主機呼叫，主機執行

device:設備呼叫，設備執行

 

執行一般c函數  funtion();

執行cuda函數   function<<>> ();

 

解釋：在GPU上面執行函數可以自定分配grid和線程，grid包含線程，因為是並列執行，因此如果內容一樣數據的生成很多是不分前后的。

 

二、運行例子的方法：

建立一個CUDA8.0 Runtim模版為基礎的工程，

或建立一個c++工程，將cpp后綴改為.cu

 

建完工程后會有一部分代碼

在主函數return 0 之前加入getchar();即可運行，關於此代碼后期可慢慢研究，這里不做講解。

源碼為：

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

#include <stdio.h>

cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size);

__global__ void addKernel(int *c, const int *a, const int *b)
{
    int i = threadIdx.x;
    c[i] = a[i] + b[i];
}

int main()
{
    const int arraySize = 5;
    const int a[arraySize] = { 1, 2, 3, 4, 5 };
    const int b[arraySize] = { 10, 20, 30, 40, 50 };
    int c[arraySize] = { 0 };

    // Add vectors in parallel.
    cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "addWithCuda failed!");
        return 1;
    }

    printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",
        c[0], c[1], c[2], c[3], c[4]);

    // cudaDeviceReset must be called before exiting in order for profiling and
    // tracing tools such as Nsight and Visual Profiler to show complete traces.
    cudaStatus = cudaDeviceReset();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceReset failed!");
        return 1;
    }
    getchar();
    return 0;
}

// Helper function for using CUDA to add vectors in parallel.
cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size)
{
    int *dev_a = 0;
    int *dev_b = 0;
    int *dev_c = 0;
    cudaError_t cudaStatus;

    // Choose which GPU to run on, change this on a multi-GPU system.
    cudaStatus = cudaSetDevice(0);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
        goto Error;
    }

    // Allocate GPU buffers for three vectors (two input, one output)    .
    cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    // Copy input vectors from host memory to GPU buffers.
    cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

    cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

    // Launch a kernel on the GPU with one thread for each element.
    addKernel<<<1, size>>>(dev_c, dev_a, dev_b);

    // Check for any errors launching the kernel
    cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
        goto Error;
    }
    
    // cudaDeviceSynchronize waits for the kernel to finish, and returns
    // any errors encountered during the launch.
    cudaStatus = cudaDeviceSynchronize();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
        goto Error;
    }

    // Copy output vector from GPU buffer to host memory.
    cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

Error:
    cudaFree(dev_c);
    cudaFree(dev_a);
    cudaFree(dev_b);
    
    return cudaStatus;
}

 

 

三、實戰代碼：

例一：第一個程序hello world

#include "cuda_runtime.h" #include "device_launch_parameters.h" #include <stdio.h> #include <Windows.h> __global__ void helloFromGPU(void) { printf("Hello World from GPU!\n"); } int main(void) { // hello from cpu
 cudaError_t cudaStatus; printf("Hello World from CPU!\n"); helloFromGPU << <1, 10 >> > (); cudaDeviceReset();//重置CUDA設備釋放程序占用的資源
    system("pause"); return 0; }

無視所有錯誤直接運行即可。

在這里  helloFromGPU << <1, 10 >> >();

表示這一函數將有十個一樣的線程，也就是這個函數總計會被執行十次。

 

 

改為helloFromGPU << <2, 10 >> >();

 

 

例二：參數的傳入

#include "cuda_runtime.h" #include "device_launch_parameters.h" #include <stdio.h> #include <Windows.h> __global__ void Add(int i, int j) { int count; count = i + j; printf("\nNum is %d\n", count); } int main(void) { Add << <1, 1 >> >(10, 15); cudaDeviceReset();//重置CUDA設備釋放程序占用的資源
    system("pause"); return 0; }

 傳入參數與一般的c沒有什么不一樣之處

例三：數據的傳入和傳出

從這里開始要用到顯存分配，在這一段中，我們的數據要從內存copy到顯存上面，然后現在又要從顯存上面copy回來

這次我們定一個減法函數，在設備上執行

__global__ void Decrease(int a, int b, int *c) { *c = a - b; }

我的要傳的數為一個整數型的c， 一般會定義一個全局的以cuda錯誤處理返回值為類型的函數，在這一函數內執行數據的傳輸，及時返回錯誤

 

cudaError_t  addWithCuda(int *c);

 

在例子中我省略了這個直接用void類型

 

void addWithCuda(int *c);

 

代碼：

#include "cuda_runtime.h" #include "device_launch_parameters.h" #include <stdio.h> #include <Windows.h>
void addWithCuda(int *c);//1.定義傳入的函數c
int main(void) { int c; c = 10; addWithCuda(&c);//2.傳入參數變量（地址）
    cudaDeviceReset();//6.重置CUDA設備釋放程序占用的資源
    printf("Value is %d", c);//7.主機上打印顯示數據
 system("pause"); return 0; } __global__ void Decrease(int a, int b, int *c) { *c = a - b; } void addWithCuda(int *c) { int *dev_c = 0;//這個相當於內存和顯存有一樣的 //3.請求CUDA設備的內存（顯存），執行CUDA函數
    cudaMalloc((void**)&dev_c, sizeof(int)); Decrease << <1, 1 >> >(15, 30, dev_c); //4.等待設備所有線程任務執行完畢
 cudaDeviceSynchronize(); //5.數據復制到主機，釋放占用空間
    cudaMemcpy(c, dev_c, sizeof(int), cudaMemcpyDeviceToHost); cudaFree(dev_c); }

 

例四：數據的傳入和傳出Ⅱ

如果要復制數據進去，那么我們先要修改下上一個例子的函數

 

1.傳入的數值全改為指針類型

__global__ void Decrease(int *a, int *b, int *c) { *c = *a - *b; }

 

2.修改函數的傳入參數

void addWithCuda(int *c,int *a,int *b);//1.定義傳入的函數c

 

3.增加處理函數中對於相應存儲空間的的申請和釋放

void addWithCuda(int *c, int *a, int *b) { int *dev_c = 0; int *dev_a = 0; int *dev_b = 0; //3.請求CUDA設備的內存（顯存），執行CUDA函數
    cudaMalloc((void**)&dev_c, sizeof(int)); cudaMalloc((void**)&dev_a, sizeof(int)); cudaMalloc((void**)&dev_b, sizeof(int)); //4.從主機復制數據到設備上
    cudaMemcpy(dev_a, a, sizeof(int), cudaMemcpyHostToDevice); cudaMemcpy(dev_b, b, sizeof(int), cudaMemcpyHostToDevice); Decrease << < 1, 1 >> >(dev_a, dev_b, dev_c); //5.等待設備所有線程任務執行完畢
 cudaDeviceSynchronize(); //6.數據復制到主機，釋放占用空間
    cudaMemcpy(c, dev_c, sizeof(int), cudaMemcpyDeviceToHost); cudaFree(dev_c); cudaFree(dev_a); cudaFree(dev_b); }

 

4.最后是主函數

int main(void) { int c; int a, b; c = 10; a = 30; b = 15; addWithCuda(&c, &a, &b);//2.傳入參數變量（地址）
    cudaDeviceReset();//7.重置CUDA設備釋放程序占用的資源
    printf("Value is %d", c);//8.主機上打印顯示數據
    system("pause"); return 0; }

 

5.代碼：

#include "cuda_runtime.h" #include "device_launch_parameters.h" #include <stdio.h> #include <Windows.h> __global__ void Decrease(int *a, int *b, int *c) { *c = *a - *b; } void addWithCuda(int *c, int *a, int *b) { int *dev_c = 0; int *dev_a = 0; int *dev_b = 0; //3.請求CUDA設備的內存（顯存），執行CUDA函數
    cudaMalloc((void**)&dev_c, sizeof(int)); cudaMalloc((void**)&dev_a, sizeof(int)); cudaMalloc((void**)&dev_b, sizeof(int)); //4.從主機復制數據到設備上
    cudaMemcpy(dev_a, a, sizeof(int), cudaMemcpyHostToDevice); cudaMemcpy(dev_b, b, sizeof(int), cudaMemcpyHostToDevice); Decrease << < 1, 1 >> >(dev_a, dev_b, dev_c); //5.等待設備所有線程任務執行完畢
 cudaDeviceSynchronize(); //6.數據復制到主機，釋放占用空間
    cudaMemcpy(c, dev_c, sizeof(int), cudaMemcpyDeviceToHost); cudaFree(dev_c); cudaFree(dev_a); cudaFree(dev_b); } int main(void) { int c; int a, b; c = 10; a = 30; b = 15; addWithCuda(&c, &a, &b);//2.傳入參數變量（地址）
    cudaDeviceReset();//7.重置CUDA設備釋放程序占用的資源
    printf("Value is %d", c);//8.主機上打印顯示數據
    system("pause"); return 0; }

 

最后再放一個程序關於循環可以自行理解

程序實現向量的加法操作，使用了一個block內部的512個線程。

#include <stdio.h>
#include<cuda_runtime.h>

//__global__聲明的函數，告訴編譯器這段代碼交由CPU調用，由GPU執行
__global__ void add(const int *dev_a,const int *dev_b,int *dev_c)
{
    int i=threadIdx.x;
    dev_c[i]=dev_a[i]+dev_b[i];
}

int main(void)
{
    //申請主機內存，並進行初始化
    int host_a[512],host_b[512],host_c[512];
    for(int i=0;i<512;i++)
    {
        host_a[i]=i;
        host_b[i]=i<<1;
    }

    //定義cudaError，默認為cudaSuccess(0)
    cudaError_t err = cudaSuccess;

    //申請GPU存儲空間
    int *dev_a,*dev_b,*dev_c;
    err=cudaMalloc((void **)&dev_a, sizeof(int)*512);
    err=cudaMalloc((void **)&dev_b, sizeof(int)*512);
    err=cudaMalloc((void **)&dev_c, sizeof(int)*512);
    if(err!=cudaSuccess)
    {
        printf("the cudaMalloc on GPU is failed");
        return 1;
    }
    printf("SUCCESS");
    //將要計算的數據使用cudaMemcpy傳送到GPU
    cudaMemcpy(dev_a,host_a,sizeof(host_a),cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b,host_b,sizeof(host_b),cudaMemcpyHostToDevice);

    //調用核函數在GPU上執行。數據較少，之使用一個Block，含有512個線程
    add<<<1,512>>>(dev_a,dev_b,dev_c);
    cudaMemcpy(&host_c,dev_c,sizeof(host_c),cudaMemcpyDeviceToHost);
    for(int i=0;i<512;i++)
        printf("host_a[%d] + host_b[%d] = %d + %d = %d\n",i,i,host_a[i],host_b[i],host_c[i]);
    cudaFree(dev_a);//釋放GPU內存
    cudaFree(dev_b);//釋放GPU內存
    cudaFree(dev_c);//釋放GPU內存
    return 0 ;
}