CUDA學習2 基礎知識和Julia示例

1.修飾符

__device__ 標記的函數從一個在器件中執行的函數呼叫，在器件中執行  

__global__ 表示該函數從一個在主機中執行的函數呼叫，在器件中執行

__host__表示在主機中呼叫，在主機中執行的函數

2.核函數

如下引用自青竹居士的博文CUDA核函數參數示意：Kernel<<<Dg,Db, Ns, S>>>(param list)。

核函數是GPU每一個thread上運行的程序。必須經過__gloabl__函數類型限定符定義。形式以下： 

                __global__ void kernel(param list){  } 

核函數只能在主機端調用，調用時必須申明執行參數。調用形式以下：

                Kernel<<<Dg,Db, Ns, S>>>(param list); 

<<<>>>運算符內是核函數的執行參數，告訴編譯器運行時如何啓動核函數，用於說明內核函數中的線程數量，以及線程是如何組織的。 

<<<>>>運算符對kernel函數完整的執行配置參數形式是<<<Dg, Db, Ns, S>>> 

• 參數Dg用於定義整個grid的維度和尺寸，即一個grid有多少個block。爲dim3類型。Dim3 Dg(Dg.x, Dg.y, 1)表示grid中每行有Dg.x個block，每列有Dg.y個block，第三維恆爲1(目前一個核函數只有一個grid)。整個grid中共有Dg.x*Dg.y個block，其中Dg.x和Dg.y最大值爲65535。

• 參數Db用於定義一個block的維度和尺寸，即一個block有多少個thread。爲dim3類型。Dim3 Db(Db.x, Db.y, Db.z)表示整個block中每行有Db.x個thread，每列有Db.y個thread，高度爲Db.z。Db.x和Db.y最大值爲512，Db.z最大值爲62。 一個block中共有Db.x*Db.y*Db.z個thread。計算能力爲1.0,1.1的硬件該乘積的最大值爲768，計算能力爲1.2,1.3的硬件支持的最大值爲1024。

• 參數Ns是一個可選參數，用於設置每一個block除了靜態分配的shared Memory之外，最多能動態分配的shared memory大小，單位爲byte。不須要動態分配時該值爲0或省略不寫。
• 參數S是一個cudaStream_t類型的可選參數，初始值爲零，表示該核函數處在哪一個流之中。

gridDim：表明線程格（grid）的尺寸。

blockIdx：表明線程塊（block）在線程格（grid）中的索引值。

blockDim：表明線程塊（block）的尺寸。

3.內存

如下引用自繞樑九日的博文CUDA編程指南閱讀筆記（二）。

在GPU上CUDA線程能夠訪問到的存儲資源有不少，每一個CUDA線程擁有獨立的本地內存（local Memory）；每個線程塊（block）都有其獨立的共享內存（shared memory），共享內存對於線程塊中的每一個線程都是可見的，它與線程塊具備相同的生存時間；同時，還有一片稱爲全局內存（global memory）的區域對全部的CUDA線程都是可訪問的。

除了上述三種存儲資源之外，CUDA還提供了兩種只讀內存空間：常量內存（constant memory）和紋理內存（texture memory），同全局內存相似，全部的CUDA線程均可以訪問它們。對於一些特殊格式的數據，紋理內存提供多種尋址模式以及數據過濾方法來操做內存。這兩類存儲資源主要用於一些特殊的內存使用場合。

一個程序啓動內核函數之後，全局內存、常量內存以及紋理內存將會一直存在直到該程序結束。下面是CUDA的內存層次圖：

 CUDA的異構編程模型假定CUDA線程都運行在一個可被看作CPU協處理器的芯片上，這就使得CUDA內核函數能夠和CPU端C程序的運行並行運行，從而加快程序的運行效率。爲了達到這個效果，CUDA程序須要管理兩大塊由DRAM構成的內存區域：CPU端能夠訪問到的主機內存（host memory）以及GPU端供CUDA內核訪問到的設備內存（device memory），設備內存主要由全局內存、常量內存以及紋理內存構成。如今，CUDA程序的運行機制便很明瞭了：CPU端代碼生成原始數據，經過CUDA運行時函數庫將這些原始數據傳輸到GPU上，在CPU端啓動CUDA內核函數進行運算，而後將運算結果從設備端傳輸到主機端，計算任務便完成了。

4.Julia

相關文件源自《CUDA by Example》源碼（csdn連接）。

源碼中julia_gpu.cu須要在

cuComplex(float a, float b) : r(a), i(b)  {}

前加上「__device__」，並將相應的dll文件（glut64.dll）拷貝到項目生成文件夾。

如下便於閱讀，將cpu_bitmap.h合併到此文件。

/*
 * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.
 *
 * NVIDIA Corporation and its licensors retain all intellectual property and 
 * proprietary rights in and to this software and related documentation. 
 * Any use, reproduction, disclosure, or distribution of this software 
 * and related documentation without an express license agreement from
 * NVIDIA Corporation is strictly prohibited.
 *
 * Please refer to the applicable NVIDIA end user license agreement (EULA) 
 * associated with this source code for terms and conditions that govern 
 * your use of this NVIDIA software.
 * 
 */


#include "../common/book.h"
#pragma comment (lib, "glut64.lib")    /* link with Win64 GLUT lib */
#include "../common/GL/glut.h"
#include "../common/GL/glext.h"

//#define GET_PROC_ADDRESS( str ) glXGetProcAddress( (const GLubyte *)str )
#define DIM 1000

struct CPUBitmap {
    unsigned char    *pixels;
    int     x, y;
    void    *dataBlock;
    void(*bitmapExit)(void*);

    CPUBitmap(int width, int height, void *d = NULL) {
        pixels = new unsigned char[width * height * 4];
        x = width;
        y = height;
        dataBlock = d;
    }

    ~CPUBitmap() {
        delete[] pixels;
    }

    unsigned char* get_ptr(void) const   { return pixels; }
    long image_size(void) const { return x * y * 4; }

    void display_and_exit(void(*e)(void*) = NULL) {
        CPUBitmap**   bitmap = get_bitmap_ptr();
        *bitmap = this;
        bitmapExit = e;
        // a bug in the Windows GLUT implementation prevents us from
        // passing zero arguments to glutInit()
        int c = 1;
        char* dummy = "";
        glutInit(&c, &dummy);
        glutInitDisplayMode(GLUT_SINGLE | GLUT_RGBA);
        glutInitWindowSize(x, y);
        glutCreateWindow("bitmap");
        glutKeyboardFunc(Key);
        glutDisplayFunc(Draw);
        glutMainLoop();
    }

    // static method used for glut callbacks
    static CPUBitmap** get_bitmap_ptr(void) {
        static CPUBitmap   *gBitmap;
        return &gBitmap;
    }

    // static method used for glut callbacks
    static void Key(unsigned char key, int x, int y) {
        switch (key) {
        case 27:
            CPUBitmap*   bitmap = *(get_bitmap_ptr());
            if (bitmap->dataBlock != NULL && bitmap->bitmapExit != NULL)
                bitmap->bitmapExit(bitmap->dataBlock);
            exit(0);
        }
    }

    // static method used for glut callbacks
    static void Draw(void) {
        CPUBitmap*   bitmap = *(get_bitmap_ptr());
        glClearColor(0.0, 0.0, 0.0, 1.0);
        glClear(GL_COLOR_BUFFER_BIT);
        glDrawPixels(bitmap->x, bitmap->y, GL_RGBA, GL_UNSIGNED_BYTE, bitmap->pixels);
        glFlush();
    }
};

struct cuComplex {
    float   r;
    float   i;
    __device__ cuComplex(float a, float b) : r(a), i(b)  {}
    __device__ float magnitude2( void ) {
        return r * r + i * i;
    }
    __device__ cuComplex operator*(const cuComplex& a) {
        return cuComplex(r*a.r - i*a.i, i*a.r + r*a.i);
    }
    __device__ cuComplex operator+(const cuComplex& a) {
        return cuComplex(r+a.r, i+a.i);
    }
};

__device__ int julia( int x, int y ) {
    const float scale = 1.5;
    float jx = scale * (float)(DIM/2 - x)/(DIM/2);
    float jy = scale * (float)(DIM/2 - y)/(DIM/2);

    cuComplex c(-0.8, 0.156);
    cuComplex a(jx, jy);

    int i = 0;
    for (i=0; i<200; i++) {
        a = a * a + c;
        if (a.magnitude2() > 1000)
            return 0;
    }

    return 1;
}

__global__ void kernel( unsigned char *ptr ) {
    // map from blockIdx to pixel position
    int x = blockIdx.x;
    int y = blockIdx.y;
    int offset = x + y * gridDim.x;

    // now calculate the value at that position
    int juliaValue = julia( x, y );
    ptr[offset*4 + 0] = 255 * juliaValue;
    ptr[offset*4 + 1] = 0;
    ptr[offset*4 + 2] = 0;
    ptr[offset*4 + 3] = 255;
}

// globals needed by the update routine
struct DataBlock {
    unsigned char   *dev_bitmap;
};

int main( void ) {
    DataBlock   data;
    CPUBitmap bitmap( DIM, DIM, &data );
    unsigned char    *dev_bitmap;

    HANDLE_ERROR( cudaMalloc( (void**)&dev_bitmap, bitmap.image_size() ) );
    data.dev_bitmap = dev_bitmap;

    dim3    grid(DIM,DIM);
    kernel<<<grid,1>>>( dev_bitmap );

    HANDLE_ERROR( cudaMemcpy( bitmap.get_ptr(), dev_bitmap,
                              bitmap.image_size(),
                              cudaMemcpyDeviceToHost ) );
                              
    HANDLE_ERROR( cudaFree( dev_bitmap ) );
                              
    bitmap.display_and_exit();
}

 運行結果以下圖。