CUDA C简介
基本操作
以下是调用GPU的基本操作代码。代码作用是将两个数相加。
其中要注意的是:
1. cudaMemcpy() 函数前两个参数传递的是地址。
2. cudaMalloc() 函数原型为:
cudaError_t cudaMalloc (void **devPtr, size_t size );
所以调用时首先将第一个参数强制转换为 (void **) 类型,再取地址 & 得到之前定义的那个一维指针地址。
代码如下:
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
__global__ void add(int a, int b, int* c){
*c = a + b;
}
int main()
{
int c[10];
int* dev_c;
cudaMalloc( (void**)&dev_c, sizeof(int) ) ;
add<<<1, 1>>>(2, 7, dev_c);
// 第一个参数和第二个参数都传的是地址.
// int c;
// cudaMemcpy( &c, dev_c, sizeof(int),cudaMemcpyDeviceToHost ) ;
// 以下代码传址传的是数组c的首地址
cudaMemcpy( c, dev_c, sizeof(int),cudaMemcpyDeviceToHost ) ;
printf("%d\n", c[0]);
cudaFree(dev_c);
return 0;
}
读取GPU的信息
用如下方式读取GPU信息:
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
int main()
{
cudaDeviceProp prop;
int count;
cudaGetDeviceCount(&count);
for (int i = 0; i < count; ++i)
{
cudaGetDeviceProperties(&prop, i);
printf(" --- General Information for device %d --- \n", i);
printf("Name: %s\n", prop.name);
printf("Compute capability: %d.%d\n", prop.major, prop.minor);
printf("Clock rate: %d\n", prop.clockRate);
printf("Device copy overlap:");
if(prop.deviceOverlap)
printf("Enable\n");
else
printf("Disable\n");
printf("Kernel execition timeout:");
if(prop.kernelExecTimeoutEnabled)
puts("Enable");
else
puts("Disable");
printf(" --- Memory Information for device %d ---\n", i);
printf("Total global mem: %ld\n", prop.totalGlobalMem);
printf("Total constant mem: %ld\n", prop.totalConstMem);
printf("Max mem pitch: %ld\n", prop.memPitch);
printf("Texture Alignment: %ld\n", prop.textureAlignment);
printf(" --- MP Information for device %d --- \n", i);
printf("Multiprocessor count %d\n", prop.multiProcessorCount);
printf("Shared mem per mp : %ld\n", prop.sharedMemPerBlock);
printf("Registers per mp: %d\n", prop.regsPerBlock);
printf("Threads in warp: %d\n", prop.warpSize);
printf("Max threads per block: %d\n", prop.maxThreadsPerBlock);
printf("Max therad dimensions: (%d, %d, %d)\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);
printf("Max grid dimensions: (%d, %d, %d)\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
puts("");
}
return 0;
}
CUDA C并行编程
向量和
使用CUDA C并行编程来计算向量与向量的和:
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
const int N = 10;
__global__ void add(int* a, int* b, int* c)
{
// blockIdx - 内置变量,保存当前执行设备代码的线程块的索引
int tid = blockIdx.x;
if (tid < N)
{
c[tid] = a[tid] + b[tid];
}
}
int main()
{
int a[N], b[N], c[N];
int *dev_a, *dev_b, *dev_c;
// GPU内存分配
cudaMalloc((void**)&dev_a, N * sizeof(int));
cudaMalloc((void**)&dev_b, N * sizeof(int));
cudaMalloc((void**)&dev_c, N * sizeof(int));
// 在CPU上给a, b 赋值
for (int i = 0; i < N; ++i)
{
a[i] = -i;
b[i] = i * i;
}
// 将数组a, b复制到GPU
cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice);
// <<<a, b>>> 第一个参数a表示设备在执行核函数时使用的并行线程块的数量
// 以下代码创建N个线程块在GPU上运行
add<<<N, 1>>>(dev_a, dev_b, dev_c);
// 将结果dev_c复制到CPU
cudaMemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost);
for (int i = 0; i < N; ++i)
{
printf("%d + %d = %d\n", a[i], b[i], c[i]);
}
cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_c);
return 0;
}
Julia集
首先要在VS和系统路径中加入glut,之后头文件导入就行了。
CPU版本:
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
//#include "cuda_gl_interop.h"
#include <gl/glut.h>
#include <stdio.h>
#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,对 GLUT (OpenGl 里面的一个工具包,包含很多函数)进行初始化,这个函数必须在其它的 GLUT使用之前调用一次。其格式比较死板,一般照抄这句glutInit(&argc, argv)就可以了*/
glutInit( &c, &dummy );
/*设置显示方式,其中 GLUT_RGBA 表示使用 RGBA 颜色,与之对应的还有GLUT_INDEX(表示使用索引颜色) ;GLUT_SINGLE 表示使用单缓冲,。与之对应的还有 GLUT_DOUBLE(使用双缓冲)。*/
glutInitDisplayMode( GLUT_SINGLE | GLUT_RGBA );
/*这个也简单,设置窗口的大小*/
glutInitWindowSize( x, y );
/*根据前面设置的信息创建窗口。参数将被作为窗口的标题。注意:窗口被创建后,并不立即显示到屏幕上。需要调用 glutMainLoop 才能看到窗口。*/
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)
{
/* 如果按键按的是Esc按键,则退出程序。*/
switch (key)
{
case 27:
CPUBitmap* bitmap = *(get_bitmap_ptr());
if (bitmap->dataBlock != NULL && bitmap->bitmapExit != NULL)
bitmap->bitmapExit( bitmap->dataBlock );
}
}
// static method used for glut callbacks
/* 画图 */
static void Draw( void )
{
CPUBitmap* bitmap = *(get_bitmap_ptr());
/*设置背景颜色*/
glClearColor( 0.0, 0.0, 0.0, 1.0 );
/*清除。GL_COLOR_BUFFER_BIT 表示清除颜色*/
glClear( GL_COLOR_BUFFER_BIT );
glDrawPixels( bitmap->x, bitmap->y, GL_RGBA, GL_UNSIGNED_BYTE, bitmap->pixels );
/*保证前面的 OpenGL 命令立即执行(而不是让它们在缓冲区中等待)。其作用跟 fflush(stdout)类似。*/
glFlush();
}
};
struct cuComplex
{
float r;
float i;
cuComplex(float a, float b) : r(a), i(b) {}
float magnitude2(void) {
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<