《GPU高性能编程CUDA实战》

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 }

自动寻找设备

将目标属性填充到cudaDeviceProp结构

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

将其传递给cudaChooseDevice()
cudaChooseDevice()返回满足条件的设备ID
将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

当矢量长度超过限制时，核函数调用会失败

解决方法：将并行线程的数量看作是处理器的数量

认定每个线程在逻辑上都可以并行执行，并且硬件可以调用这些线程以便实际执行。通过将并行化过程与硬件的实际执行过程解耦开来。

步骤：

计算每个并行线程的初始化索引，以及递增的线程
对线程索引和线程块索引进行线性化，使每个并行线程从不同的索引开始

起始索引：
```
 int tid = threadIdx.x + blockIdx.c * blockDim.x;
```
对索引进行递增，递增步长为线程格中正在运行的线程数量。此数值等于每个线程块中的线程数量乘以线程格中线程块的数量，即
```
 tid += blockDim.x * gridDim.x;
```
线程块数量确定没明确说明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 }

posted on 2019-09-24 15:18 Skipig 阅读(2003) 评论(0) 编辑收藏举报

会员力量，点亮园子希望

刷新页面返回顶部

导航

公告