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类型的可选参数，初始值为零，表示该核函数处在哪个流之中。
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(); }
运行结果如下图。
