碼上快樂

相關內容    簡體    繁體

《GPU高性能編程CUDA實戰》

本文轉載自   查看原文   2019-09-24 15:18   753 

1.第一個CUDA程序

 1 #include <iostream>
 2 
 3 __global__ void kernel(void) {  //__global__告知編譯器函數kernel用設備代碼編輯器
 4 }
 5 
 6 int main() {                    //默認主機編譯
 7     kernel << <1, 1 >> > ();
 8     printf("HelloWorld");
 9     return 0;
10 }

CUDA提供與C在語言級別上集成，在主機代碼中調用設備代碼

尖括號內參數用來確定運行時如何啟動設備代碼

2.關鍵詞

 1 #include <iostream>
 2 #include <cuda_runtime.h>
 3 #include <device_launch_parameters.h>
 4 //#include <book.h>
 5 
 6 __global__ void add(int a, int b, int *c) {
 7     *c = a + b;
 8 }
 9 
10 int main() {
11     int c;
12     int *dev_c;
13     cudaMalloc((void**)&dev_c, sizeof(int));
14 
15     add <<<1, 1 >>> (2, 7, dev_c);
16 
17     cudaMemcpy(&c,dev_c,sizeof(int),cudaMemcpyDeviceToHost);
18     printf("2 + 7=%d\\n", c);
19 
20     cudaFree(dev_c);
21     return 0;
22 }

 

  • 像調用C函數一樣將參數傳遞給核函數
  • 設備執行操作時需要分配內存

使用cudaMalloc()分配內存

作用：使CUDA在運行時在設備上分配內存

cudaMalloc((void**)&dev_c, sizeof(int)

  • 第一個參數是指針，用來保存新分配內存地址變量
  • 第二個參數是分配內存的大小
  • 返回類型為void*

不能在主機代碼中對cudaMalloc()返回的指針進行解引用(Dereference)。

主機代碼可以將這個指針作為參數傳遞，對其進行算術運算，轉換為另一種不同類型，但是不可以使用這個指針來進行讀取或者寫入內存

設備指針使用限制：

  • 可以將cudaMalloc()分配的指針傳遞給設備上執行的函數
  • 可以將cudaMalloc()分配的指針傳遞給在主機上執行的函數
  • 可以在設備代碼中使用cudaMalloc()分配的指針進行內存讀/寫
  • 不能在主機代碼中使用cudaMalloc()分配的指針進行內存讀/寫

不能用標准C的free()釋放cudaMalloc()分配的內存，需要調用cudaFree()

主機上不能對設備上的內存做任何修改

訪問設備內存兩種方法

  • 在設備代碼中使用設備指針
    • 主機指針只能訪問主機代碼中的內存
    • 設備指針只能訪問設備代碼中的內存
  • 主機調用cudaMemcpy()

cudaMemcpy()

類似標准C中的memcpy()，多了一個指定設備內存指針(源指針/目標指針)的參數

void *memcpy(void *dest, const void *src, size_t n); 由src指向地址為起始地址的連續n個字節的數據復制到以destin指向地址為起始地址的空間內。 #include<string.h> 函數返回一個指向dest的指針。

  • cudaMemcpyDeviceToHost
  • cudaMemcpyHostToDevice
  • cudaMemcpyDeviceToDevice告訴運行時兩個指針均位於設備上
  • 若源指針和目標指針均位於主機上，可以直接調用memcpy()函數

3.查詢設備信息

調用cudaGetDeviceCount,返回結構參數如圖：

代碼

 1 #include <iostream>
 2 #include <cuda_runtime.h>
 3 #include <device_launch_parameters.h>
 4 int main() {
 5     cudaDeviceProp prop;
 6 
 7     int count;
 8     cudaGetDeviceCount(&count);
 9     for (int i = 0;i < count;i++) {
10         cudaGetDeviceProperties(&prop, i);
11         
12         printf(" ---General Information for Device %d---\\n", i);
13         printf("Name: %s\\n", prop.name);
14         printf("Compute capability: %d.%d\\n", prop.major,prop.minor);
15         printf("Clock rate: %d\\n", prop.clockRate);
16         printf("Device copy overlap:  ");
17         if (prop.deviceOverlap)
18             printf("Enabled\\n");
19         else
20             printf("Disabled\\n");
21         printf("Kernel execition timeout :  ");
22         if (prop.kernelExecTimeoutEnabled)
23             printf("enabled\\n");
24         else
25             printf("Disabled\\n");
26         printf("\\n");
27 
28         printf("---Memory Information for device %d---\\n", i);
29         printf("Total global Mem:%ld\\n", prop.totalGlobalMem);
30         printf("Total constant Mem:%ld\\n", prop.totalConstMem);
31         printf("Max mem pitch:%ld\\n", prop.memPitch);
32         printf("Texture Alignment:%ld\\n", prop.textureAlignment);
33         printf("\\n");
34 
35         printf("---MP Information for device %d---\\n", i);
36         printf("Multiprocessor count  :%d\\n", prop.multiProcessorCount);
37         printf("Shared mem per mp:%ld\\n", prop.sharedMemPerBlock);
38         printf("Registers per mp: %d\\n", prop.regsPerBlock);
39         printf("Threads in warp: %d\\n", prop.warpSize);
40         printf("Max threads per block: %d\\n", prop.maxThreadsPerBlock);
41         printf("Max thread dimensions:(%d, %d, %d)\\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);
42         printf("Max grid dimensions:(%d, %d, %d)\\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
43         printf("\\n");
44 
45     }
46 }

 

自動尋找設備

  1. 將目標屬性填充到cudaDeviceProp結構

     cudaDeviceProp prop;
 memset(&prop,0,sizeof(cudaDeviceProp));
 prop.major=1;
 prop.minor=3;

  2. 將其傳遞給cudaChooseDevice()

  3. cudaChooseDevice()返回滿足條件的設備ID

  4. 將ID傳遞給cudaSetDevice()，之后所有操作在此設備上進行

完整程序

 1 #include <iostream>
 2 #include <cuda_runtime.h>
 3 #include <device_launch_parameters.h>
 4 int main() {
 5     cudaDeviceProp prop;
 6     int dev;
 7 
 8     cudaGetDevice(&dev);
 9     printf("ID of current CUDA device: %d\\n", dev);
10 
11     memset(&prop, 0, sizeof(cudaDeviceProp));
12     prop.major = 1;
13     prop.minor = 3;
14     cudaChooseDevice(&dev, &prop);
15     printf("ID of CUDA device closest to reviaion 1.3: %d\\n", dev);
16     cudaSetDevice(dev);
17 }

 

設備使用

速度快->多核處理器的GPU

核函數與CPU有密集交互->在集成的GPU上運行代碼，因為其可與CPU共享內存

NVIDIA的SLI（Scalable Link Interface，可伸縮鏈路接口）技術使得多個獨立的GPU可以並排排列。

無論是哪種情況，應用程序都可以從多個GPU中選擇最適合的GPU。

如果應用程序依賴於GPU的某些特定屬性，或者需要在系統中最快的GPU上運行，此API有幫助，因為CUDA運行時本身並不能保證為應用程序選擇最優或者最合適的GPU。

小結

CUDA C/C++只是對標准C/C++進行了語言級擴展，利用修改符指定代碼在主機或設備上運行。

__global__指明函數在GPU上運行

使用GPU上內存，通過與C相關API對應的CUDA的API

對比

4.CUDA C並行編程

GPU計算應用前景取決於能否從問題中發掘出大規模並行性

書籍P29，對CPU上並行進行了否定

 1 #include <iostream>
 2 #include <cuda_runtime.h>
 3 #include <device_launch_parameters.h>
 4 
 5 #define N 10000
 6 
 7 __global__ void add(int *a, int *b, int *c) {
 8     int tid = blockIdx.x; //計算位於此索引處的數據
 9         if (tid < N)
10             c[tid] = a[tid] + b[tid];
11 }
12 
13 int main() {
14     int a[N], b[N], c[N];
15     int *dev_a, *dev_b, *dev_c;
16 
17     //GPU上分配內存
18     cudaMalloc((void**)&dev_a, N * sizeof(int));
19     cudaMalloc((void**)&dev_b, N * sizeof(int));
20     cudaMalloc((void**)&dev_c, N * sizeof(int));
21 
22     //對數組a,b賦值
23     for (int i = 0;i < N;i++) {
24         a[i] = -i;
25         b[i] = i*i;
26     }
27 
28     //HostToDevice
29     cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
30     cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice);
31 
32     add << <N, 1 >> > (dev_a, dev_b, dev_c);
33 
34     //將結果從GPU復制到CPU
35     cudaMemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost);
36 
37     //輸出結果
38     for (int i = 0;i < N;i++) {
39         printf("%d + %d = %d\\n", a[i], b[i], c[i]);
40     }
41     //釋放內存
42     cudaFree(dev_a);
43     cudaFree(dev_b);
44     cudaFree(dev_c);
45 
46     return 0;
47 }

 

上例僅給出函數main()，其在GPU上的實現與在CPU上的實現是不同的，但此時無差別

kernel<<<N，1>>>(dev_a, dev_b, dev_c);

第一個參數表示設備在執行核函數時使用的並行線程塊數量，運行N個核函數副本,前行線程塊集合也稱為一個線程格grid

  • 在核函數中，通過變量blockIdx.x確定當前運行區塊
  • blockIdx.x為當前執行設備代碼的線程塊的索引

e.g.N=4，此時4個線程的的blockIdx.x值分別為0，1，2，3

每個線程塊實際執行的代碼如下：

4.1實例

Julia集：通過迭代等式對復平面中的等求值。

  • 迭代等式計算結果發散，朝無窮大的方向增長，此點不屬於Julia集合
  • 迭代等式收斂，位於某個邊界滿園之內，此點屬於Julia集合

迭代等式：

$$Z_{n+1}^2=Z_{n}^2+C$$

4.1.1基於CPU的Julia集

 1 #include <stdio.h>
 2 
 3 #include <cuda_runtime.h>
 4 #include <device_launch_parameters.h>
 5 
 6 #include "D:\\common\\book.h"
 7 #include "D:\\common\\cpu_bitmap.h"
 8 
 9 #define DIM 1000
10 
11 //計算在復數上進行，定義結構保存復數
12 //定義復數的加法和乘法運算
13 struct cuComplex {
14     float r;//實部r
15     float i;//虛部i
16     cuComplex(float a,float b):r(a),i(b){}
17     float magnitude2() { return r / r + i + i; }
18     cuComplex operator*(const cuComplex &a) {
19         return cuComplex(r*a.r - i*a.i, i*a.r + r*a.i);
20     }
21     cuComplex operator*(const cuComplex &a) {
22         return cuComplex(r + a.r, i + a.i);
23     }
24 };
25 
26 int julia(int x, int y) {
27     //實現圖形綻放的scale因數
28     const float scale = 1.5;
29 
30     //將像素坐標轉換為空間坐標
31     //像素移動DIM/2個單位，將復平面原點定位在圖像中心
32     //圖像范圍在-1.0到1.0，圖像坐標綻放了DIM/2倍
33     float jx = scale*(float)(DIM / 2 - x) / (DIM / 2);
34     float jy = scale*(float)(DIM / 2 - y) / (DIM / 2);
35 
36     //迭代公式中的C為-0.5+0.156i
37     cuComplex c(-0.8, 0.156);
38     cuComplex a(jx, jy);
39 
40     int i = 0;
41     for (i = 0;i < 200;i++) {
42         //a = a*a + c;
43         if (a.magnitude2() > 1000)//迭代結果閾值
44             return 0;
45     }
46     return 1;
47 }
48 
49 
50 
51 
52 //核函數對繪制的所有點進行迭代
53 void kernel(unsigned char *ptr) {
54     for (int y = 0;y < DIM;y++) {
55         for (int x = 0;x < DIM;x++) {
56             int offset = x + y*DIM;
57 
58             //調用julia()判斷點是否屬於Julia集
59             //是返回1，點為紅色
60             //否返回0，點為黑色，可改
61             int juliaValue = julia(x, y);
62             ptr[ offset * 4 + 0 ] = 255 * juliaValue;
63             ptr[ offset * 4 + 1 ] = 0;
64             ptr[ offset * 4 + 2 ] = 0;
65             ptr[ offset * 4 + 31 ] = 255;
66         }
67     }
68 }
69 
70 int main() {
71     CPUBitmap bitmap(DIM, DIM); //通過工具庫創建位圖圖像
72     unsigned char *ptr = bitmap.get_ptr;
73 
74     //將指向位圖數據的指針傳遞給核函數
75     kernel(ptr);
76 
77     bitmap.display_and_exit();
78 
79     return 0;
80 }

 

4.1.2基於GPU的Julia集

 1 #include <stdio.h>
 2 #include <cuda_runtime.h>
 3 #include <device_launch_parameters.h>
 4 
 5 #include "D:\\common\\book.h"
 6 #include "D:\\common\\cpu_bitmap.h"
 7 
 8 #define DIM 1000
 9 
10 //計算在復數上進行，定義結構保存復數
11 //定義復數的加法和乘法運算
12 struct cuComplex {
13     float r;//實部r
14     float i;//虛部i
15     __device__ cuComplex(float a, float b) :r(a), i(b) {}
16     __device__ float magnitude2() {
17         return r * r + i * i;
18     }
19     __device__ cuComplex operator*(const cuComplex &a) {
20         return cuComplex(r*a.r - i*a.i, i*a.r + r*a.i);
21     }
22     __device__ cuComplex operator+(const cuComplex &a) {
23         return cuComplex(r + a.r, i + a.i);
24     }
25 };
26 
27 
28 //判斷點是否屬於Julia集
29 __device__ int julia(int x, int y) {
30     const float scale = 1.5;
31     float jx = scale*(float)(DIM / 2 - x) / (DIM / 2);
32     float jy = scale*(float)(DIM / 2 - y) / (DIM / 2);
33 
34     cuComplex c(-0.8, 0.156);
35     cuComplex a(jx, jy);
36 
37     int i = 0;
38     for (i = 0;i < 200;i++) {
39         a = a*a + c;
40         if (a.magnitude2() > 1000)
41             return 0;
42     }
43     return 1;
44 }
45 
46 
47 
48 
49 //不需要for()來生成像素索引傳遞給julia()
50 //cuda運行時在變量blockIdx中包含這些索引
51 //在聲明線程格時，線程格每一維的大小與圖像每一維的大小是相等的，因此
52 //在（0,1）到（DIM,DIM）之間每個像素點都能分配一個線程塊
53 
54 __global__ void kernel(unsigned char *ptr) {
55     //將threadIdx/BlockIdx映射到像素位置
56     int x = blockIdx.x;
57     int y = blockIdx.y;
58     //內置變量gridDim，常數，保存線程格每一維大小
59     //行索引乘以線程格寬度+列索引得到ptr唯一索引，范圍(DIM*DIM-1)
60     int offset = x + y*gridDim.x;
61 
62     //計算此位置上的值
63     int juliaValue = julia(x, y);
64     ptr[offset * 4 + 0] = 255 * juliaValue;
65     ptr[offset * 4 + 1] = 0;
66     ptr[offset * 4 + 2] = 0;
67     ptr[offset * 4 + 3] = 255;
68 }
69 
70 int main() {
71     //創建DIM*DIM大小的位圖圖像
72     CPUBitmap bitmap(DIM, DIM);
73     //保存設備上數據的副本
74     unsigned char *dev_bitmap;
75 
76     cudaMalloc((void**)&dev_bitmap, bitmap.image_size());
77 
78     dim3 grid(DIM, DIM);
79     kernel << <grid, 1 >> >(dev_bitmap);
80 
81     //返回計算結果
82     cudaMemcpy(bitmap.get_ptr(), dev_bitmap, bitmap.image_size(), cudaMemcpyDeviceToHost);
83 
84     bitmap.display_and_exit();
85     cudaFree(dev_bitmap);
86 }

 

計算結果

 

計算線程塊需要的數據索引

  • 核函數的每個副本可以通過內置變量blockIdx來判斷哪個線程塊在執行它
  • 通過內置變量gridDim獲得線程格的大小

5線程協作

kernel<<<N,1>>>
  • 第一個參數是啟動的線程塊數量
  • CUDA運行時每個線程塊中創建的線程數量
    • 啟動的總線程數量 N個線程塊*1個線程/線程塊=N個並行線程

5.1矢量求和

5.1.1使用線程實現GPU上矢量求和

改動：

  • add<<<N,1>>>(dev_a,dev_b,dev_c) -> add<<<1,N>>>(dev_a,dev_b,dev_c)

  • 數據索引方法線程塊索引變為線程索引

      int tid = blockIdx.x; -> int tid = threadIdx.x;

完整程序

 1 #include <iostream>
 2 #include <cuda_runtime.h>
 3 #include <device_launch_parameters.h>
 4 
 5 #define N 10000
 6 
 7 __global__ void add(int *a, int *b, int *c) {
 8     int tid = threadIdx.x; //計算位於此索引處的數據
 9         if (tid < N)
10             c[tid] = a[tid] + b[tid];
11 }
12 
13 int main() {
14     int a[N], b[N], c[N];
15     int *dev_a, *dev_b, *dev_c;
16 
17     //GPU上分配內存
18     cudaMalloc((void**)&dev_a, N * sizeof(int));
19     cudaMalloc((void**)&dev_b, N * sizeof(int));
20     cudaMalloc((void**)&dev_c, N * sizeof(int));
21 
22     //對數組a,b賦值
23     for (int i = 0;i < N;i++) {
24         a[i] = -i;
25         b[i] = i*i;
26     }
27 
28     //HostToDevice
29     cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
30     cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice);
31 
32     add << <1, N >> > (dev_a, dev_b, dev_c);
33 
34     //將結果從GPU復制到CPU
35     cudaMemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost);
36 
37     //輸出結果
38     for (int i = 0;i < N;i++) {
39         printf("%d + %d = %d\\n", a[i], b[i], c[i]);
40     }
41     //釋放內存
42     cudaFree(dev_a);
43     cudaFree(dev_b);
44     cudaFree(dev_c);
45 
46     return 0;
47 }

 

GPU上對更長矢量求和

  • 線程塊每一維的數量限制為65535
  • 啟動核函數時每個線程塊中的線程數量不能超過設備屬性結構中maxThreadsPerBlock域的值 大部分是每個線程塊512個線程

更改

核函數中的索引計算方法 核函數的調用方式

計算索引方法類似於將二維索引空間轉換為線性空間的標准算法

int tid =threadIdx.x + blockIdx.x * blockDim.x

gridDim 線程格中每一維的線程塊數量 二維

blockDim 線程塊中每一維的線程數量 三維

int offset = x + y * DIM;

DIM表示線程塊大小即線程的數量

y為線程塊索引，x為線程塊中的線程索引

計算得到索引：tid = threadIdx.x + blockIdx.x * blockDim.x

核函數調用

kernel <<<(N+127/128,128)>>>(dev_a，dev_b，dev_c)

啟動128個線程

N+127/128一種向上取整的算法，計算大於或等於N的128的最小倍數

對於多啟動的線程，在訪問輸入數組和輸出數組之前，檢查線程的偏移是否位於0到N之間

if(tid<N)    c[tid] = a[tid] + b[tid];

當索引越過數組邊界時，核函數將自動停止計算，核函數不對越過數組邊界的內存進行讀取或寫入

GPU上對任意長度的矢量求和

線程塊每一維的數量限制為65535

當矢量長度超過限制時，核函數調用會失敗

解決方法：將並行線程的數量看作是處理器的數量

認定每個線程在邏輯上都可以並行執行，並且硬件可以調用這些線程以便實際執行。通過將並行化過程與硬件的實際執行過程解耦開來。

步驟：

  1. 計算每個並行線程的初始化索引，以及遞增的線程

  2. 對線程索引和線程塊索引進行線性化，使每個並行線程從不同的索引開始

    起始索引：

     int tid = threadIdx.x + blockIdx.c * blockDim.x;

  3. 對索引進行遞增，遞增步長為線程格中正在運行的線程數量。此數值等於每個線程塊中的線程數量乘以線程格中線程塊的數量，即

     tid += blockDim.x * gridDim.x;

  4. 線程塊數量確定沒明確說明P59

     add<<<128,128>>>(dev_a,dev_b,dev_c);

總的程序：

 1 #include <stdio.h>
 2 #include <cuda_runtime.h>
 3 #include <device_launch_parameters.h>
 4 
 5 #include "D:\\common\\book.h"
 6 //#include "D:\\common\\cpu_bitmap.h"
 7 
 8 #define N (33*1024)
 9 
10 __global__ void add(int *a, int *b, int *c) {
11     int tid = threadIdx.x + blockIdx.x*blockDim.x;
12     while (tid < N) {
13         c[tid] = a[tid] + b[tid];
14         tid += blockDim.x*gridDim.x;
15     }
16 }
17 
18 int main() {
19     int a[N], b[N], c[N];
20     int *dev_a, *dev_b, *dev_c;
21 
22     //GPU上分配內存
23     cudaMalloc((void**)&dev_a, N * sizeof(int));
24     cudaMalloc((void**)&dev_b, N * sizeof(int));
25     cudaMalloc((void**)&dev_c, N * sizeof(int));
26 
27     //CPU上為數組a,b賦值
28     for (int i = 0;i < N;i++) {
29         a[i] = i;
30         b[i] = i*i;
31     }
32 
33     //將數組a,b復制到GPU
34     cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
35     cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice);
36     add<<<128,128>>>(dev_a, dev_b, dev_c);
37 
38     //將數組c復制回CPU
39     cudaMemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost);
40 
41     //驗證GPU完成了工作
42     bool success = true;
43     for (int i = 0;i < N;i++) {
44         if ((a[i] + b[i]) != c[i]) {
45             printf("Error: %d + %d != %d\\n", a[i], b[i], c[i]);
46             success = false;
47         }
48     }
49     if (success)
50         printf("done\\n");
51 
52     //釋放GPU上內存
53     cudaFree(dev_a);
54     cudaFree(dev_b);
55     cudaFree(dev_c);
56 
57     return 0;
58 }

 

5.2.2在GPU上使用線程實現波紋效果

 1 #include "D:/common/book.h"
 2 #include "D:/common/cpu_anim.h"
 3 
 4 #define DIM 1024
 5 
 6 struct DataBlock {
 7     unsigned char *dev_bitmap;
 8     CPUAnimBitmap *bitmap;
 9 };
10 
11 void cleanup(DataBlock *d) {
12     cudaFree(d->dev_bitmap);
13 }
14 
15 __global__ void kernel(unsigned char* ptr, int ticks) {
16     //將threadIdx、BlockIdx映射到像素位置
17     //線程得到其在線程塊中的索引，及此線程塊在線程格中的索引，並將兩值轉換為圖形的唯一索引(x,y)
18     int x = threadIdx.x + blockIdx.x * blockDim.x;
19     int y = threadIdx.y + blockIdx.y * blockDim.y;
20 
21     //對x,y進行線性化得到輸出緩沖區中的一個偏移
22     int offset = x + y * blockDim.x * gridDim.x;
23     //int offset = y + x * blockDim.y * gridDim.y;//這兩個offset等效
24 
25     float fx = x - DIM / 2;
26     float fy = y - DIM / 2;
27     float d = sqrtf(fx * fx + fy * fy);
28     unsigned char grey = (unsigned char)(128.0f + 127.0f * cos(d / 10.0f - ticks / 7.0f) / (d / 10.0f + 1.0f));
29     ptr[offset * 4 + 0] = grey;//grey 2D時間函數
30     ptr[offset * 4 + 1] = grey;
31     ptr[offset * 4 + 2] = grey;
32     ptr[offset * 4 + 3] = 255;
33 }
34 
35 void generate_frame(DataBlock *d, int ticks) {
36     dim3 blocks(DIM / 16, DIM / 16);    //聲明一個二維變量，線程格中包含的並行線程塊數量
37     dim3 threads(16, 16);                //聲明一個二維變量，線程塊中包含的線程數量
38     
39     //核函數來計算像素值
40     //指針指向保存輸出像素值的設備內存，是全局變量，其指向的內存是在main()中華西的。全局性針對主機，參數要傳遞讓設備能夠訪問到
41     //當前時效ticks傳遞給generate_frame(),核函數根據當前動畫時間生成正確的幀
42     kernel << <blocks, threads >> > (d->dev_bitmap, ticks);
43     HANDLE_ERROR(cudaMemcpy(d->bitmap->get_ptr(),
44         d->dev_bitmap,
45         d->bitmap->image_size(),
46         cudaMemcpyDeviceToHost));
47 }
48 
49 int main() {
50     DataBlock data;
51     CPUAnimBitmap bitmap(DIM, DIM, &data); //大部分復雜性隱藏在輔助類CPUAnimBitmap中
52     data.bitmap = &bitmap;
53     HANDLE_ERROR(cudaMalloc((void**)&data.dev_bitmap, bitmap.image_size()));
54 
55     //將指向generate_frame()函數的指針傳遞給anim_and_exit()，每當生成一幀新的動畫，都將調用generate_frame()
56     bitmap.anim_and_exit((void(*)(void*, int))generate_frame, (void(*)(void*))cleanup);
57 }

 


免責聲明！

本站轉載的文章為個人學習借鑒使用，本站對版權不負任何法律責任。如果侵犯了您的隱私權益，請聯系本站郵箱yoyou2525@163.com刪除。



猜您在找 GPGPU OpenCL/CUDA 高性能編程的10大注意事項 高性能JavaScript DOM編程 CUDA Pro：通過向量化內存訪問提高性能 低價租用高性能GPU進行深度學習 高性能緩存 Caffeine 原理及實戰 《CUDA並行程序設計：GPU編程指南》 Linux高性能服務器編程 conda 安裝GPU——CUDA gpu命令cuda命令 顯卡、GPU和CUDA簡介
 
粵ICP備18138465號   © 2018-2025 CODEPRJ.COM