CHG: CUDA

顯示具有 CUDA 標籤的文章。顯示所有文章

2018年5月31日星期四

ubuntu 16.04 安裝 cuda 9.0

驅動程式可以不安裝過程會自己安裝，先安裝也會重新安裝cuda選用的版本，整個流程跑完要手動重新開機。

因為有安裝驅動，建議下面流程跑完就馬上重新開機，不然用一用可能會當機。安裝驅動之後重新開機第一次會比較久，等一下會黑畫面什麼都沒有很像死機了。

過程涉及安裝驅動~如果有發生什麼衝突可能導致系統崩潰無法開啟，這點要注意一下~

下載位置：CUDA Toolkit 9.0 Downloads

可以不用點有包含再底下指令（cuda版本有升級文章還沒更新就要手動改版本號）

安裝指令

# 加入驅動鏡像站
sudo add-apt-repository ppa:graphics-drivers/ppa -y

# 升級系統
sudo apt-get update
sudo apt-get upgrade -y

# 安裝 cuda
cd ~/Downloads
sudo wget -c https://developer.nvidia.com/compute/cuda/9.0/Prod/local_installers/cuda-repo-ubuntu1604-9-0-local_9.0.176-1_amd64-deb
sudo dpkg -i cuda-repo-ubuntu1604-9-0-local_9.0.176-1_amd64-deb
sudo apt-key add /var/cuda-repo-9-0-local/7fa2af80.pub
sudo apt-get update
sudo apt-get install cuda -y

# 加入函式庫
export PATH=/usr/local/cuda-9.0/bin${PATH:+:${PATH}}

到這邊就裝好了，還要配置他的bin跟lib位置，否則直接打 nvcc 會出現找不到
方法是參考官方說明網站：docs.nvidia.com

# 函式庫標頭檔路徑
export PATH=/usr/local/cuda-9.0/bin${PATH:+:${PATH}}

# 函式庫路徑(這個好像不用)
export LD_LIBRARY_PATH=/usr/local/cuda-9.0/lib64 ${LD_LIBRARY_PATH:+:${LD_LIBRARY_PATH}}

這樣就配置好了可以打

nvcc --version

只不過剛剛的配置方法只有對這個終端機有用，換終端機或是重開就沒了
想要真正加入的話把那兩行寫入 vi .bashrc 這個文件裏面開機自動執行就好了

vi .bashrc

加入最尾端或是隨意找個地方寫也可以，然後重新登入再打測試版本的指令就有了。
最後測試一下到底有沒有裝對。

這邊記得先重開機一次否則測試結果會失敗歐

git clone https://github.com/hunandy14/cuda_copyArr
cd cuda_copyArr
make run

看一下有沒有跑出 test ok 有的話就是一切都正常了，包含驅動。（失敗的話重新開機再試試看）

這是一個簡單的cuda程式，其中我有寫一個比較容易的呼叫使用記憶體的函式，適合新手學習使用。

寫文章測試的時候 cuda9.0 選用的版本 384 而當前最新版本是 390 ，建議不要更動這個版本裝好就別升級了，可能會發生一些無法預測的錯誤。

2018年5月30日星期三

ubuntu 16.04 安裝 NVIDIA 驅動快捷方法

安裝驅動

為了避免未知的相依性問題建議先更新系統到最新

sudo apt-get update
sudo apt-get upgrade

然後我們新增驅動的庫

sudo add-apt-repository ppa:graphics-drivers/ppa
sudo apt-get update

然後搜尋驅動管理員打開，就有最新版驅動可以安裝了

software-properties-gtk --open-tab=4

選擇新版的安裝~

這邊選擇390是因為我到 NVIDIA 官方去下載我的顯示卡 GT960 最新到390板所以選擇這個版本，建議去官方查一下比較保險~
http://www.nvidia.com/Download/index.aspx?lang=en-uk

裝好重新開機一下，然後打指令檢查

lspci -vnn | grep -i VGA -A 12

最後 Kernel modules: 可以看到你的驅動程式版本

參考

https://standbymesss.blogspot.com/2016/09/ubuntu1404-nvidiadriver.html

2018年1月12日星期五

cuda bilinear 紋理記憶體、共享記憶體實作與解說

之前介紹過一篇

[原始碼] C / C++ 線性插補 bilinear 與注意事項

裡面詳細介紹注意事項與如何實作，大致簡單說一下就是等比取值

如果 A=0, B=10 ，那麼 AB=1 大概就是這個原理，如果是二維平面4個點就是取3次剛剛的過程，先兩次兩個點，然後再把兩個點的結果取一次。

CPU函式

上面那個敘述就是我們的主要函式了，大致可以寫成如下程式

__host__ __device__
inline static float bilinearRead(const float* img, 
    size_t width, float y, float x) // 線性取值
{
    // 獲取鄰點(不能用 1+)
    size_t x0 = floor(x);
    size_t x1 = ceil(x);
    size_t y0 = floor(y);
    size_t y1 = ceil(y);
    // 獲取比例(只能用 1-)
    float dx1 = x - x0;
    float dx2 = 1 - dx1;
    float dy1 = y - y0;
    float dy2 = 1 - dy1;
    // 獲取點
    const float& A = img[y0*width + x0];
    const float& B = img[y0*width + x1];
    const float& C = img[y1*width + x0];
    const float& D = img[y1*width + x1];
    // 乘出比例(要交叉)
    float AB = A*dx2 + B*dx1;
    float CD = C*dx2 + D*dx1;
    float X = AB*dy2 + CD*dy1;
    return X;
}

要操作這個函式大致上就用兩個for用x, y跑過圖中所有的點就可以了

__host__ void biliner_CPU_core(vector<float>& img, const vector<float>& img_ori, 
    size_t width, size_t height, float Ratio)
{
    int newH = static_cast<int>(floor(height * Ratio));
    int newW = static_cast<int>(floor(width  * Ratio));
    img.resize(newH*newW);
    // 跑新圖座標
    for (int j = 0; j < newH; ++j) {
        for (int i = 0; i < newW; ++i) {
            // 調整對齊
            float srcY, srcX;
            if (Ratio < 1) {
                srcY = ((j+0.5f)/Ratio) - 0.5;
                srcX = ((i+0.5f)/Ratio) - 0.5;
            } else {
                srcY = j * (height-1.f) / (newH-1.f);
                srcX = i * (width -1.f) / (newW-1.f);
            }
            // 獲取插補值
            img[j*newW + i] = bilinearRead(img_ori.data(), width, srcY, srcX);
        }
    }
}

共享記憶體

共享記憶體這邊有引用到之前寫的函式庫，詳細可以參考這篇站內文
https://charlottehong.blogspot.tw/2018/01/cuda.html

// 共享記憶體線性取值核心
__global__ void biliner_share_kernel(float* dst, const float* src, int srcW, int srcH, float ratio) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    int newH = (int)(floor(srcH * ratio));
    int newW = (int)(floor(srcW * ratio));
    if(i < srcW*ratio && j < srcH*ratio) { // 會多跑一點點要擋掉
        // 調整對齊
        float srcY, srcX;
        if (ratio < 1) {
            srcY = ((j+0.5f)/ratio) - 0.5;
            srcX = ((i+0.5f)/ratio) - 0.5;
        } else {
            srcY = j * (srcH-1.f) / (newH-1.f);
            srcX = i * (srcW -1.f) / (newW-1.f);
        }
        // 獲取插補值
        dst[j*newW + i] = bilinearRead(src, srcW, srcY, srcX);
    }
}

差不多就是從CPU那個函式改過來的而已，使用時候要先new兩個空間

// 共享記憶體線性取值函式
__host__ void biliner_share_core(float *dst, const float* src,
    size_t srcW, size_t srcH, float ratio)
{
    Timer T; T.priSta = 1;
    // 設置GPU所需長度
    int srcSize = srcW*srcH;
    int dstSize = srcSize*ratio*ratio;

    // 要求GPU空間
    T.start();
    CudaData<float> gpu_src(srcSize);
    T.print("  GPU new 空間1");
    T.start();
    CudaData<float> gpu_dst(dstSize);
    T.print("  GPU new 空間2");
    // 複製到GPU
    T.start();
    gpu_src.memcpyIn(src, srcSize);
    T.print("  GPU 複製");

    // 設置執行緒
    dim3 block(BLOCK_DIM, BLOCK_DIM);
    dim3 grid(ceil((float)srcW*ratio / BLOCK_DIM), ceil((float)srcH*ratio / BLOCK_DIM));
    T.start();
    biliner_share_kernel <<< grid, block >> > (gpu_dst, gpu_src, srcW, srcH, ratio);
    T.print("  核心計算");

    // 取出GPU值
    T.start();
    gpu_dst.memcpyOut(dst, dstSize);
    T.print("  GPU 取出資料");

    // 釋放GPU空間
    T.start();
    gpu_src.~CudaData();
    gpu_dst.~CudaData();
    T.print("  GPU 釋放空間");
}

紋理記憶體

這個速度最快代碼也最少，可以直接調用cuda內優化過的方法

// 宣告GPU紋理變數(只能放全域)
texture<float, 2, cudaReadModeElementType> rT;
// 紋理線性取值核心
__global__ void biliner_texture_kernel(float* dst, int srcW, int srcH, float ratio) {
    int idxX = blockIdx.x * blockDim.x + threadIdx.x,
        idxY = blockIdx.y * blockDim.y + threadIdx.y;
    if(idxX < srcW*ratio && idxY < srcH*ratio) { // 會多跑一點點要擋掉
        float srcX = idxX / ratio;
        float srcY = idxY / ratio;
        size_t idx = (idxY*srcW*ratio + idxX);
        dst[idx] = tex2D(rT, srcX+0.5, srcY+0.5);
    }
}

使用

__host__ void biliner_texture_core(float *dst, const float* src,
    size_t dstW, size_t dstH, float ratio)
{
    Timer T; T.priSta = 1;
    // 設置GPU所需長度
    int srcSize = dstW*dstH;
    int dstSize = srcSize*ratio*ratio;

    // 宣告 texture2D陣列並綁定
    T.start();
    CudaMemArr<float> cuArray(src, dstW, dstH);
    cudaBindTextureToArray(rT, cuArray);
    T.print("  GPU new 紋理空間+複製");

    // 設置 插植模式and超出邊界補邊界
    rT.filterMode = cudaFilterModeLinear;
    rT.addressMode[0] = cudaAddressModeClamp;
    rT.addressMode[1] = cudaAddressModeClamp;

    // 要求GPU空間
    T.start();
    CudaData<float> gpu_dst(dstSize);
    T.print("  GPU new 一般空間");

    // 設置執行緒
    dim3 block(BLOCK_DIM, BLOCK_DIM);
    dim3 grid(ceil((float)dstW*ratio / BLOCK_DIM), ceil((float)dstH*ratio / BLOCK_DIM));
    T.start();
    biliner_texture_kernel <<< grid, block >> > (gpu_dst, dstW, dstH, ratio);
    T.print("  核心計算");

    // 取出GPU值
    T.start();
    gpu_dst.memcpyOut(dst, dstSize);
    T.print("  GPU 取出資料");
}

總結

可以看出大部分時間其實是消耗再第一次new空間的時候，至於記憶體的複製貌似也沒消耗太多時間，建議圖可以直接讀進共享記憶體就好了，不需要先用主機記憶體讀然後再複製進去，從一開始的時候就宣告共享記憶體。

texture times = 0.327s
share times = 0.4s
CPU times = 2.578s

範例用的原始碼：https://github.com/hunandy14/Cuda_ImgBilinear

2018年1月8日星期一

cuda 自動管理記憶體函式庫

因為還沒有實際在專案上跑過，可能有bug或是考慮不周的情況發生，再自行修改~

另外目前最新版 VS2017 還不支援，有兩個辦法可以處理

使用VS2015
安裝舊版VS2017並不更新
在最新版2017新增2015的套件，讓方案用2015跑

3看起來是最好的，不過實際運行有一些bug，每次編譯完之後更改程式碼在按除錯並不會編譯新的程式碼會直接執行舊的exe，清除重編可以不過每次都按會很麻煩。最佳手段是選1或2。

詳情參考：https://charlottehong.blogspot.tw/2017/12/visutal-studio-2017-cuda90.html

cuda 通常流程

宣告空間

新建 共享記憶體 這一個cpu跟gpu都可以用，一般都是用這個傳入gpu

T* gpuData;
cudaMalloc((void**)&gpuData, size*sizeof(T));

輸入資料

cudaMemcpy(gpuData, dataIn, size*sizeof(T), cudaMemcpyHostToDevice);

輸出資料

cudaMemcpy(dataIn, gpuData, size*sizeof(T), cudaMemcpyDeviceToHost);

cuda 記憶體控制類別

上述流程繁雜，不過基本上就是跟C語言一樣，為了寫作方便可以讓一個class來管理記憶體，下面是我寫的類別。

稍微有點長我把它放到外部空間gist

使用

基本用法都包含進去，需要什麼在自己包什麼。
下面是使用範例

/*****************************************************************
Name : 
Date : 2018/01/08
By   : CharlotteHonG
Final: 2018/01/08
*****************************************************************/
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <iostream>
#include <vector>
using namespace std;

#include "timer.hpp"
#include "CudaData.cuh"

__global__ void cudacopy(float* b, float* a, int size){
    // 乾式寫法
    const int idx = blockIdx.x*blockDim.x + threadIdx.x;
    if(idx<size){
        b[idx]=a[idx];
    }
    // 迴圈寫法
    //for(int i=threadIdx.x; i<size; i+=blockDim.x){
        //b[i]=a[i];
    //}
}
void cpucopoy(float* b, float* a, int size) {
    for(int i=0; i<size; ++i){
        b[i]=a[i];
    }
}
void testCuda(size_t size) {
    Timer T;
    // 配置主機記憶體
    vector<float> img_data(size), cpu_data(size), gpu_data(size);
    float* a = img_data.data(); // 原始資料
    float* b = cpu_data.data(); // CPU計算後資料
    float* c = gpu_data.data(); // GPU輸出回來資料
    // 設置初值
    float test_val=7;
    for(int i=0; i<size; i++){
        a[i]=test_val;
    }
    // 配置顯示記憶體, 載入資料.
    T.start();
    CudaData<float> gpuDataIn(a, size), gpuDataOut(size);
    gpuDataOut.memset(0, size);
    T.print("  Cuda Data malloc and copy");

    // 網格區塊設定. (與 kernel for 的次數有關)
    const size_t blkDim=16;
    int grid(size/blkDim+1);  // 網格要含蓋所有範圍, 所以除完要加 1.
    int block(blkDim);        // 區塊設定 16x16.
    // Cuda Kernel 執行運算
    T.start();
    cudacopy<<<grid,block>>>(gpuDataOut, gpuDataIn, size);
    T.print("  Cuda-copy");
    // 取出GPU資料
    gpuDataOut.memcpyOut(c, size);

    // CPU 執行運算
    T.start();
    cpucopoy(b, a, size);
    T.print("  Cpu-copy");
    // 測試
    bool f=0;
    for(size_t i = 0; i < size; i++) {
        if(c[i] != b[i]) {f=1;}
    }
    // 測試報告
    if(f==0) {
        cout << "test ok" << endl;
    } else {
        cout << "test Error" << endl;
    }
}
int main(){
    Timer T;
    testCuda(99999);
    T.print("ALL time.");
    return 0;
}

2017年12月26日星期二

Cuda 9.0 範例程式使用 for 迴圈複製陣列

現在9版的安裝 cuda 好像沒有什麼難度了，順對弄對先安裝好 VisualStudio 然後再安裝 cuda 就一路下一步一直按到底就好了。

然後開啟VC就有專案可以開啟了

範例程序有for迴圈可以參考，這裡我把它簡化了一下一些流程並加上時間監控

首先先是計時函式庫

/*****************************************************************
Name : Timer.hpp
Date : 2017/12/19
By   : CharlotteHonG
Final: 2017/12/26
*****************************************************************/
#pragma once
#include <iostream>
#include <string>
#include <ctime>

class Timer {
public:
    Timer(std::string name=""): name(name){
        startTime = clock();
    }
    operator double() {
        return time;
    }
public:
    void start() {
        startTime = clock();
        flag=0;
    }
    void end() {
        finalTime = clock();
        time = (double)(finalTime - startTime)/CLOCKS_PER_SEC;
    }
    void print(std::string name="") {
        if (flag==0) {
            flag = 1;
            end();
        }
        if(name=="") {
            name=this->name;
        }
        if(priSta) {
            std::cout << "#" << name << ", " << " time = " << time << "s" << std::endl;
        }
    }
private:
    std::string name;
    clock_t startTime;
    clock_t finalTime;
    double time;
    bool flag = 0;
public:
    bool priSta = 1;
};

再來是控管 cuda 記憶體的類別，他可以用來自動要求空間並複製進去以及在程序結束時自動解構，不需要再 free() cuda記憶體。

由於這個有寫樣板所以沒法簡單把實作拆開，建議就直接寫在cuh裡面省事吧~

/*****************************************************************
Name : cudaData.cuh
Date : 2017/12/19
By   : CharlotteHonG
Final: 2017/12/19
*****************************************************************/
// Cuda 記憶體自動管理程序
template <class T>
class CudaData {
public:
    CudaData(){}
    CudaData(size_t size){
        malloc(size);
    }
    CudaData(T* dataIn ,size_t size): len(size){
        memcpyInAuto(dataIn, size);
    }
    ~CudaData(){
        if(gpuData!=nullptr) {
            cudaFree(gpuData);
            gpuData = nullptr;
            len = 0;
        }
    }
public:
    void malloc(size_t size) {
        this->~CudaData();
        len = size;
        cudaMalloc((void**)&gpuData, size*sizeof(T));
    }
    void memcpyIn(T* dataIn ,size_t size) {
        if(size > len) {throw out_of_range("memcpyIn input size > curr size.");}
        cudaMemcpy(gpuData, dataIn, size*sizeof(T), cudaMemcpyHostToDevice);
    }
    void memcpyInAuto(T* dataIn ,size_t size) {
        malloc(size);
        memcpyIn(dataIn, size);
    }
    void memcpyOut(T* dataIn ,size_t size) {
        cudaMemcpy(dataIn, gpuData, size*sizeof(T), cudaMemcpyDeviceToHost);
    }
    void memset(int value, size_t size) {
        if(size>len) {
            throw out_of_range("memset input size > curr size.");
        }
        cudaMemset(gpuData, value, size*sizeof(T));
    }
    size_t size() {
        return this->len;
    }
public:
    operator T*() {
        return gpuData;
    }
private:
    T* gpuData;
    size_t len=0;
};

主程式

再來是主程式了，會用到上面兩個函式庫


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

#include <iostream>
#include <vector>
using namespace std;

#include "timer.hpp"
#include "CudaData.cuh"

__global__ void cudacopy(float* b, float* a, int size){
    // 乾式寫法
    const int idx = blockIdx.x*blockDim.x + threadIdx.x;
    if(idx<size){
        b[idx]=a[idx];
    }
    // 迴圈寫法
    for(int i=threadIdx.x; i<size; i+=blockDim.x){
        b[i]=a[i];
    }
}
void cpucopoy(float* b, float* a, int size) {
    for(int i=0; i<size; ++i){
        b[i]=a[i];
    }
}
void testCuda(size_t size) {
    Timer T;

    // 配置主機記憶體
    vector<float> img_data(size), cpu_data(size), gpu_data(size);
    float* a = img_data.data(); // 原始資料
    float* b = cpu_data.data(); // CPU計算後資料
    float* c = gpu_data.data(); // GPU輸出回來資料
    // 設置初值
    float test_val=7;
    for(int i=0; i<size; i++){
        a[i]=test_val;
    }

    // 配置顯示記憶體, 載入資料.
    T.start();
    CudaData<float> gpuDataIn(a, size), gpuDataOut(size);
    gpuDataOut.memset(0, size);
    T.print("  Cuda Data malloc and copy");

    // 網格區塊設定. (與 kernel for 的次數有關)
    const size_t blkDim=16;
    int grid(size/blkDim+1);  // 網格要含蓋所有範圍, 所以除完要加 1.
    int block(blkDim);        // 區塊設定 16x16.

    // Cuda Kernel 執行運算
    T.start();
    cudacopy<<<grid,block>>>(gpuDataOut, gpuDataIn, size);
    T.print("  Cuda-copy");
    // 取出GPU資料
    gpuDataOut.memcpyOut(c, size);

    // CPU 執行運算
    T.start();
    cpucopoy(b, a, size);
    T.print("  Cpu-copy");

    // 測試
    bool f=0;
    for(size_t i = 0; i < size; i++) {
        if(c[i] != b[i]) {f=1;}
    }

    // 測試報告
    if(f==0) {
        cout << "test ok" << endl;
    } else {
        cout << "test Error" << endl;
    }
}

int main(){
    Timer T;
    testCuda(1000);
    T.print("ALL time.");
    return 0;
}

訂閱：文章 (Atom)

2018年5月31日 星期四

ubuntu 16.04 安裝 cuda 9.0

ubuntu 16.04 安裝 cuda 9.0

2018年5月30日 星期三

ubuntu 16.04 安裝 NVIDIA 驅動 快捷方法

ubuntu 16.04 安裝 NVIDIA 驅動 快捷方法

安裝驅動

參考

2018年1月12日 星期五

cuda bilinear 紋理記憶體、共享記憶體實作與解說

cuda bilinear 紋理記憶體、共享記憶體實作與解說

CPU函式

共享記憶體

紋理記憶體

總結

2018年1月8日 星期一

cuda 自動管理記憶體 函式庫

cuda 自動管理記憶體 函式庫

cuda 通常流程

宣告空間

輸入資料

輸出資料

cuda 記憶體控制類別

使用

2017年12月26日 星期二

Cuda 9.0 範例程式 使用 for 迴圈複製陣列

Cuda 9.0 範例程式 使用 for 迴圈複製陣列

主程式

2018年5月31日星期四

2018年5月30日星期三

ubuntu 16.04 安裝 NVIDIA 驅動快捷方法

ubuntu 16.04 安裝 NVIDIA 驅動快捷方法

2018年1月12日星期五

2018年1月8日星期一

cuda 自動管理記憶體函式庫

cuda 自動管理記憶體函式庫

2017年12月26日星期二

Cuda 9.0 範例程式使用 for 迴圈複製陣列

Cuda 9.0 範例程式使用 for 迴圈複製陣列