璇﹁В CUDA By Example 涓殑 Julia Set 缁樺埗GPU浼樺寲

Posted 2021-03-12

鏍囩锛?a href='http://www.mamicode.com/so/1/spl' title='spl'>spl   bit   soft   grid   error   tin   瀹屾垚   dia   缁樺埗   

绗旇€呮祴璇曠幆澧僔S2019銆?/p>

鍩烘湰浠嬬粛

鍘熶功浣滆€呭紩鍏ulia Sets鎰忓湪浣跨敤GPU鍔犻€熷浘褰㈢殑缁樺埗銆侸ulia Set 鏄寚婊¤冻涓嬪紡杩唬鏀舵暃鐨勫鏁伴泦鍚?br /> [ Z_{n+1}=Z_{n}^2+C ]

鐜閰嶇疆

璺戣繖涓緥瀛愮殑涓昏鍥伴毦搴旇鍦ㄤ簬閰嶇疆鐜銆傝繖涓▼搴忎緷璧栦簬openGL涓殑glut搴撱€傜敱浜嶸S2019鐨勬暣涓蒋浠舵灦鏋勫彂鐢熶簡寰堝ぇ鍙樺寲锛屼竴浜涢摼鎺ュ簱鍜屽ご鏂囦欢鐨勪綅缃兘鍙戠敓浜嗘敼鍙橈紝鍥犳涓€浜涙枃绔犱腑鐨勯厤缃柟娉曞け鏁堜簡銆?/p>

棣栧厛鎴戜滑闇€瑕佽幏鍙杇lut搴撶殑澶存枃浠朵互鍙婂姩鎬侀摼鎺ュ簱銆?/p>

鐐瑰嚮杩欓噷cg-toolkit鑾峰彇銆傚畨瑁呮垚鍔熶箣鍚庯紝鎵惧埌C:Program Files (x86)NVIDIA CorporationCg銆傛敞鎰忓嬀閫夊畨瑁呴€夐」鐨剎64鐩稿叧搴旂敤銆?/p>

灏嗗叾涓殑lib鏂囦欢澶逛腑鐨刜glut32.lib_澶嶅埗鍒癈:Program Files (x86)Windows Kits10Lib10.0.18362.0ucrtx86

灏嗗叾涓殑lib.x64鏂囦欢澶逛腑鐨?em>glut32.lib澶嶅埗鍒癈:Program Files (x86)Windows Kits10Lib10.0.18362.0ucrtx64骞朵笖閲嶅懡鍚嶅叾涓?em>glut64.lib

绗旇€呰繍琛岀殑鏄?4浣嶇郴缁燂紝灏卞皢bin.x64涓殑_glut32.dll_澶嶅埗鍒癈:WindowsSystem32涓?/p>

鍦ㄨ繖閲?a href="https://cloud.tsinghua.edu.cn/d/d1c3964cd86a46cc9623/">涓嬭浇澶存枃浠?/a>銆備笅杞藉畬鎴愪箣鍚庯紝灏嗗ご鏂囦欢鎷疯礉鍒癈:Program Files (x86)Windows Kits10Include10.0.18362.0ucrt銆傚苟寤虹珛鏂囦欢澶笹L鎶婂畠浠寘鎷捣鏉ャ€?/p>

鎻愮ず锛屾牳蹇冩槸鎵惧埌C:Program Files (x86)Windows Kits10锛屼笉瑕佸湪Microsoft Visual Studio鏂囦欢澶归噷娴垂鏃堕棿銆?/p>

鍚庨潰鐨?0.0.18362.0鏍规嵁鐗堟湰涓嶅悓鍙兘涓嶄竴鑷达紝鍏蜂綋闂鍏蜂綋鍒嗘瀽

杩欎釜浠ｇ爜杩橀渶瑕佷竴浜涘埆鐨勫ご鏂囦欢銆傚gl_helper.h, book.h, cpu_bitmap.h 绛?鍦ㄨ繖閲?a href="https://cloud.tsinghua.edu.cn/d/8aaf938f871149dc83e5/">涓嬭浇鍚庡鍒跺埌C:Program FilesNVIDIA GPU Computing ToolkitCUDAv10.2include

涓昏浠ｇ爜

CPU Julia Set

RGBA妯″紡涓紝姣忎竴涓儚绱犱細淇濆瓨浠ヤ笅鏁版嵁锛歊鍊硷紙绾㈣壊鍒嗛噺锛夈€丟鍊硷紙缁胯壊鍒嗛噺锛夈€丅鍊硷紙钃濊壊鍒嗛噺锛夊拰A鍊硷紙alpha鍒嗛噺锛夈€傚叾涓孩銆佺豢銆佽摑涓夌棰滆壊鐩哥粍鍚堬紝灏卞彲浠ュ緱鍒版垜浠墍闇€瑕佺殑鍚勭棰滆壊锛岃€宎lpha涓嶇洿鎺ュ奖鍝嶉鑹诧紝瀹冪殑鍚箟鏄€忔槑搴︺€?a href="https://www.cnblogs.com/tjulym/p/5037124.html%22openGL鐨勯鑹?22">1

涓嬮潰鏄函绮笴PU涓殑浠ｇ爜锛屽熀鏈殑娉ㄩ噴鍦ㄤ唬鐮佷腑

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "device_functions.h"
#include "device_atomic_functions.h"
#include <cuda.h>
#include "book.h"
#include <cpu_bitmap.h>
#include <stdio.h>

#define DIM 1000        //鍥惧儚鐨勫儚绱犺竟闀垮ぇ灏?
struct cuComplex
{
    float r;
    float i;
    cuComplex(float a, float b) : r(a), i(b) {}
    float magnitude2() { return r * r + i * i; }    //璁＄畻澶嶆暟鐨勬ā鍊?    cuComplex operator* (const cuComplex& a)
    {
        return cuComplex(r * a.r - i * a.i, i * a.r + r * a.i);
    }
    cuComplex operator+ (const cuComplex& a)
    {
        return cuComplex(r + a.r, i + a.i);
    }
};

int julia(int x, int y)
{
    const float scale = 1.5;    //鏀惧ぇ鍊嶇巼
    float jx = scale * (float)(DIM / 2 - x) / (DIM / 2);    //鍧愭爣鍙樻崲锛屾姇褰卞埌-1~1scale
    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;
}

void kernel(unsigned char* ptr)
{
    for (int y = 0; y < DIM; y++)   //閬嶅巻鏁翠釜bitmap
    {
        for (int x = 0; x < DIM; x++)
        {
            int offset = x + y * DIM;
            int juliaValue = julia(x, y);
            //娉ㄦ剰openGL杩欓噷鐨勯鑹叉牸寮忔槸RGBA,000涓洪粦鑹?            ptr[offset * 4 + 0] = 255 * juliaValue;
            ptr[offset * 4 + 1] = 0;
            ptr[offset * 4 + 2] = 0;
            ptr[offset * 4 + 3] = 255;
        }
    }
}

int main()
{
    CPUBitmap bitmap(DIM, DIM);
    unsigned char* ptr = bitmap.get_ptr();
    kernel(ptr);    //杩愯娓叉煋
    bitmap.display_and_exit();
}

GPU Julia Set

娉ㄦ剰鐢变簬鍐呮牳鍑芥暟鏄?kbd>global鐨勶紝瑕佸湪GPU涓婅繍琛岄渶瑕佸皢鍏惰皟鐢ㄧ殑julia鍑芥暟鍔犱笂device銆傚張鍥犱负锛宒evice鍑芥暟鍙兘鐢眃evice鍑芥暟鎴栬€単lobal鍑芥暟璋冪敤锛屾墍浠ユ渶濂芥妸缁撴瀯浣撲腑鐨勬墍鏈夊嚱鏁伴兘鍔犱笂device銆?/p>

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "device_functions.h"
#include "device_atomic_functions.h"
#include <cuda.h>
#include "book.h"
#include <cpu_bitmap.h>
#include <stdio.h>

//灏忎簬65536
#define DIM 1000        //鍥惧儚鐨勫儚绱犺竟闀垮ぇ灏?
struct cuComplex
{
    float r;
    float i;
    __device__ cuComplex(float a, float b) : r(a), i(b) {}
    __device__ float magnitude2() { 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);    //鍧愭爣鍙樻崲锛屾姇褰卞埌-1~1scale
    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)
{
    int x = blockIdx.x; //绾靛悜绾跨▼绱㈠紩(x鏂瑰悜鏈濆彸锛屾槸琛?
    int y = blockIdx.y; //绾靛悜绾跨▼绱㈠紩(y鏂瑰悜鏈濅笅锛屾槸鍒楋級
    int offset = x + y * gridDim.x;
    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;
}


int main()
{
    CPUBitmap bitmap(DIM, DIM);
    unsigned char* dev_bitmap;
    //鍦℅PU涓垎閰嶇┖闂?    HANDLE_ERROR(cudaMalloc((void**)&dev_bitmap, bitmap.image_size()));
    dim3 grid(DIM, DIM);    //dim3缁撴瀯浣?    kernel <<<grid, 1 >>> (dev_bitmap); //涓€涓嚎绋嬪潡涓殑绾跨▼缃戠粶1000x1000
    HANDLE_ERROR(cudaMemcpy(bitmap.get_ptr(), dev_bitmap, bitmap.image_size(), cudaMemcpyDeviceToHost));    //灏哾ev_bitmap涓殑鍐呭浠巇evice鎷疯礉鍒癱pu涓?    bitmap.display_and_exit();
    HANDLE_ERROR(cudaFree(dev_bitmap));
}

---

鍙傝€冭祫鏂?/p>