程序博客网 > 淘宝上买被子靠谱吗

CUDA编程札记

来源：互联网 发布：淘宝上买被子靠谱吗 编辑：程序博客网 时间：2024/05/22 00:33






[cpp] view plaincopy
  1. const int N = 33 * 1024;  
  2. const int threadsPerBlock = 256;  
  3. const int blocksPerGrid =  
  4.             imin( 32, (N+threadsPerBlock-1) / threadsPerBlock );  
  5.   
  6.   
  7. __global__ void dot( float *a, float *b, float *c ) {  
  8.     __shared__ float cache[threadsPerBlock];  
  9.     int tid = threadIdx.x + blockIdx.x * blockDim.x;  
  10.     int cacheIndex = threadIdx.x;  
  11.   
  12.     float   temp = 0;  
  13.     while (tid < N) {  
  14.         temp += a[tid] * b[tid];  
  15.         tid += blockDim.x * gridDim.x;  
  16.     }  
  17.       
  18.     // set the cache values  
  19.     cache[cacheIndex] = temp;  
  20.       
  21.     // synchronize threads in this block  
  22.     __syncthreads();  
  23.   
  24.     // for reductions, threadsPerBlock must be a power of 2  
  25.     // because of the following code  
  26.     int i = blockDim.x/2;  
  27.     while (i != 0) {  
  28.         if (cacheIndex < i)  
  29.             cache[cacheIndex] += cache[cacheIndex + i];  
  30.         __syncthreads();  
  31.         i /= 2;  
  32.     }  
  33.   
  34.     if (cacheIndex == 0)  
  35.         c[blockIdx.x] = cache[0];  
  36. }  
  37.   
  38.   
  39. int main( void ) {  
  40.     float   *a, *b, c, *partial_c;  
  41.     float   *dev_a, *dev_b, *dev_partial_c;  
  42.   
  43.     // allocate memory on the cpu side  
  44.     a = (float*)malloc( N*sizeof(float) );  
  45.     b = (float*)malloc( N*sizeof(float) );  
  46.     partial_c = (float*)malloc( blocksPerGrid*sizeof(float) );  
  47.   
  48.     // allocate the memory on the GPU  
  49.     HANDLE_ERROR( cudaMalloc( (void**)&dev_a,  
  50.                               N*sizeof(float) ) );  
  51.     HANDLE_ERROR( cudaMalloc( (void**)&dev_b,  
  52.                               N*sizeof(float) ) );  
  53.     HANDLE_ERROR( cudaMalloc( (void**)&dev_partial_c,  
  54.                               blocksPerGrid*sizeof(float) ) );  
  55.   
  56.     // fill in the host memory with data  
  57.     for (int i=0; i<N; i++) {  
  58.         a[i] = i;  
  59.         b[i] = i*2;  
  60.     }  
  61.   
  62.     // copy the arrays 'a' and 'b' to the GPU  
  63.     HANDLE_ERROR( cudaMemcpy( dev_a, a, N*sizeof(float),  
  64.                               cudaMemcpyHostToDevice ) );  
  65.     HANDLE_ERROR( cudaMemcpy( dev_b, b, N*sizeof(float),  
  66.                               cudaMemcpyHostToDevice ) );   
  67.   
  68.     dot<<<blocksPerGrid,threadsPerBlock>>>( dev_a, dev_b,  
  69.                                             dev_partial_c );  
  70.   
  71.     // copy the array 'c' back from the GPU to the CPU  
  72.     HANDLE_ERROR( cudaMemcpy( partial_c, dev_partial_c,  
  73.                               blocksPerGrid*sizeof(float),  
  74.                               cudaMemcpyDeviceToHost ) );  
  75.   
  76.     // finish up on the CPU side  
  77.     c = 0;  
  78.     for (int i=0; i<blocksPerGrid; i++) {  
  79.         c += partial_c[i];  
  80.     }  
  81.   
  82.     #define sum_squares(x)  (x*(x+1)*(2*x+1)/6)  
  83.     printf( "Does GPU value %.6g = %.6g?\n", c,  
  84.              2 * sum_squares( (float)(N - 1) ) );  
  85.   
  86.     // free memory on the gpu side  
  87.     HANDLE_ERROR( cudaFree( dev_a ) );  
  88.     HANDLE_ERROR( cudaFree( dev_b ) );  
  89.     HANDLE_ERROR( cudaFree( dev_partial_c ) );  
  90.   
  91.     // free memory on the cpu side  
  92.     free( a );  
  93.     free( b );  
  94.     free( partial_c );  
  95. }  


[cpp] view plaincopy
  1. struct Lock {  
  2.     int *mutex;  
  3.     Lock( void ) {  
  4.         HANDLE_ERROR( cudaMalloc( (void**)&mutex,sizeof(int) ) );  
  5.         HANDLE_ERROR( cudaMemset( mutex, 0, sizeof(int) ) );  
  6.     }  
  7.     ~Lock( void ) {  
  8.         cudaFree( mutex );  
  9.     }  
  10.     __device__ void lock( void ) {  
  11.         while( atomicCAS( mutex, 0, 1 ) != 0 );  
  12.     }  
  13.     __device__ void unlock( void ) {  
  14.         atomicExch( mutex, 0 );  
  15.     }  
  16. };  

[cpp] view plaincopy
  1. #define imin(a,b) (a<b?a:b)  
  2.   
  3. const int N = 33 * 1024 * 1024;  
  4. const int threadsPerBlock = 256;  
  5. const int blocksPerGrid =  
  6.             imin( 32, (N+threadsPerBlock-1) / threadsPerBlock );  
  7.   
  8. __global__ void dot( Lock lock, float *a,  
  9.                      float *b, float *c ) {  
  10.     __shared__ float cache[threadsPerBlock];  
  11.     int tid = threadIdx.x + blockIdx.x * blockDim.x;  
  12.     int cacheIndex = threadIdx.x;  
  13.   
  14.     float   temp = 0;  
  15.     while (tid < N) {  
  16.         temp += a[tid] * b[tid];  
  17.         tid += blockDim.x * gridDim.x;  
  18.     }  
  19.       
  20.     // set the cache values  
  21.     cache[cacheIndex] = temp;  
  22.       
  23.     // synchronize threads in this block  
  24.     __syncthreads();  
  25.   
  26.     // for reductions, threadsPerBlock must be a power of 2  
  27.     // because of the following code  
  28.     int i = blockDim.x/2;  
  29.     while (i != 0) {  
  30.         if (cacheIndex < i)  
  31.             cache[cacheIndex] += cache[cacheIndex + i];  
  32.         __syncthreads();  
  33.         i /= 2;  
  34.     }  
  35.   
  36.     if (cacheIndex == 0) {  
  37.         // wait until we get the lock  
  38.         lock.lock();  
  39.        // we have the lock at this point, update and release  
  40.         *c += cache[0];  
  41.         lock.unlock();  
  42.     }  
  43. }  
  44.   
  45.   
  46. int main( void ) {  
  47.     float   *a, *b, c = 0;  
  48.     float   *dev_a, *dev_b, *dev_c;  
  49.   
  50.     // allocate memory on the cpu side  
  51.     a = (float*)malloc( N*sizeof(float) );  
  52.     b = (float*)malloc( N*sizeof(float) );  
  53.   
  54.     // allocate the memory on the GPU  
  55.     HANDLE_ERROR( cudaMalloc( (void**)&dev_a,  
  56.                               N*sizeof(float) ) );  
  57.     HANDLE_ERROR( cudaMalloc( (void**)&dev_b,  
  58.                               N*sizeof(float) ) );  
  59.     HANDLE_ERROR( cudaMalloc( (void**)&dev_c,  
  60.                               sizeof(float) ) );  
  61.   
  62.     // fill in the host memory with data  
  63.     for (int i=0; i<N; i++) {  
  64.         a[i] = i;  
  65.         b[i] = i*2;  
  66.     }  
  67.   
  68.     // copy the arrays 'a' and 'b' to the GPU  
  69.     HANDLE_ERROR( cudaMemcpy( dev_a, a, N*sizeof(float),  
  70.                               cudaMemcpyHostToDevice ) );  
  71.     HANDLE_ERROR( cudaMemcpy( dev_b, b, N*sizeof(float),  
  72.                               cudaMemcpyHostToDevice ) );   
  73.     HANDLE_ERROR( cudaMemcpy( dev_c, &c, sizeof(float),  
  74.                               cudaMemcpyHostToDevice ) );   
  75.   
  76.     Lock    lock;  
  77.     dot<<<blocksPerGrid,threadsPerBlock>>>( lock, dev_a,  
  78.                                             dev_b, dev_c );  
  79.   
  80.     // copy c back from the GPU to the CPU  
  81.     HANDLE_ERROR( cudaMemcpy( &c, dev_c,  
  82.                               sizeof(float),  
  83.                               cudaMemcpyDeviceToHost ) );  
  84.   
  85.     #define sum_squares(x)  (x*(x+1)*(2*x+1)/6)  
  86.     printf( "Does GPU value %.6g = %.6g?\n", c,  
  87.              2 * sum_squares( (float)(N - 1) ) );  
  88.   
  89.     // free memory on the gpu side  
  90.     HANDLE_ERROR( cudaFree( dev_a ) );  
  91.     HANDLE_ERROR( cudaFree( dev_b ) );  
  92.     HANDLE_ERROR( cudaFree( dev_c ) );  
  93.   
  94.     // free memory on the cpu side  
  95.     free( a );  
  96.     free( b );  
  97. }  

[cpp] view plaincopy
  1. __global__ void histo_kernel( unsigned char *buffer,  
  2.                               long size,  
  3.                               unsigned int *histo ) {  
  4.     // calculate the starting index and the offset to the next  
  5.     // block that each thread will be processing  
  6.     int i = threadIdx.x + blockIdx.x * blockDim.x;  
  7.     int stride = blockDim.x * gridDim.x;  
  8.     while (i < size) {  
  9.         atomicAdd( &histo[buffer[i]], 1 );  
  10.         i += stride;  
  11.     }  
  12. }  
  13.   
  14. int main( void ) {  
  15.     unsigned char *buffer =  
  16.                      (unsigned char*)big_random_block( SIZE );  
  17.   
  18.     // capture the start time  
  19.     // starting the timer here so that we include the cost of  
  20.     // all of the operations on the GPU.  
  21.     cudaEvent_t     start, stop;  
  22.     HANDLE_ERROR( cudaEventCreate( &start ) );  
  23.     HANDLE_ERROR( cudaEventCreate( &stop ) );  
  24.     HANDLE_ERROR( cudaEventRecord( start, 0 ) );  
  25.   
  26.     // allocate memory on the GPU for the file's data  
  27.     unsigned char *dev_buffer;  
  28.     unsigned int *dev_histo;  
  29.     HANDLE_ERROR( cudaMalloc( (void**)&dev_buffer, SIZE ) );  
  30.     HANDLE_ERROR( cudaMemcpy( dev_buffer, buffer, SIZE,  
  31.                               cudaMemcpyHostToDevice ) );  
  32.   
  33.     HANDLE_ERROR( cudaMalloc( (void**)&dev_histo,  
  34.                               256 * sizeof( int ) ) );  
  35.     HANDLE_ERROR( cudaMemset( dev_histo, 0,  
  36.                               256 * sizeof( int ) ) );  
  37.   
  38.     // kernel launch - 2x the number of mps gave best timing  
  39.     cudaDeviceProp  prop;  
  40.     HANDLE_ERROR( cudaGetDeviceProperties( &prop, 0 ) );  
  41.     int blocks = prop.multiProcessorCount;  
  42.     histo_kernel<<<blocks*2,256>>>( dev_buffer, SIZE, dev_histo );  
  43.       
  44.     unsigned int    histo[256];  
  45.     HANDLE_ERROR( cudaMemcpy( histo, dev_histo,  
  46.                               256 * sizeof( int ),  
  47.                               cudaMemcpyDeviceToHost ) );  
  48.   
  49.     // get stop time, and display the timing results  
  50.     HANDLE_ERROR( cudaEventRecord( stop, 0 ) );  
  51.     HANDLE_ERROR( cudaEventSynchronize( stop ) );  
  52.     float   elapsedTime;  
  53.     HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime,  
  54.                                         start, stop ) );  
  55.     printf( "Time to generate:  %3.1f ms\n", elapsedTime );  
  56.   
  57.     long histoCount = 0;  
  58.     for (int i=0; i<256; i++) {  
  59.         histoCount += histo[i];  
  60.     }  
  61.     printf( "Histogram Sum:  %ld\n", histoCount );  
  62.   
  63.     // verify that we have the same counts via CPU  
  64.     for (int i=0; i<SIZE; i++)  
  65.         histo[buffer[i]]--;  
  66.     for (int i=0; i<256; i++) {  
  67.         if (histo[i] != 0)  
  68.             printf( "Failure at %d!  Off by %d\n", i, histo[i] );  
  69.     }  
  70.   
  71.     HANDLE_ERROR( cudaEventDestroy( start ) );  
  72.     HANDLE_ERROR( cudaEventDestroy( stop ) );  
  73.     cudaFree( dev_histo );  
  74.     cudaFree( dev_buffer );  
  75.     free( buffer );  
  76.     return 0;  
  77. }  


[cpp] view plaincopy
  1. __global__ void histo_kernel( unsigned char *buffer,  
  2.                               long size,  
  3.                               unsigned int *histo ) {  
  4.   
  5.     // clear out the accumulation buffer called temp  
  6.     // since we are launched with 256 threads, it is easy  
  7.     // to clear that memory with one write per thread  
  8.     __shared__  unsigned int temp[256];  
  9.     temp[threadIdx.x] = 0;  
  10.     __syncthreads();  
  11.   
  12.     // calculate the starting index and the offset to the next  
  13.     // block that each thread will be processing  
  14.     int i = threadIdx.x + blockIdx.x * blockDim.x;  
  15.     int stride = blockDim.x * gridDim.x;  
  16.     while (i < size) {  
  17.         atomicAdd( &temp[buffer[i]], 1 );  
  18.         i += stride;  
  19.     }  
  20.     // sync the data from the above writes to shared memory  
  21.     // then add the shared memory values to the values from  
  22.     // the other thread blocks using global memory  
  23.     // atomic adds  
  24.     // same as before, since we have 256 threads, updating the  
  25.     // global histogram is just one write per thread!  
  26.     __syncthreads();  
  27.     atomicAdd( &(histo[threadIdx.x]), temp[threadIdx.x] );  
  28. }  
  29.   
  30. int main( void ) {  
  31.     unsigned char *buffer =  
  32.                      (unsigned char*)big_random_block( SIZE );  
  33.   
  34.     // capture the start time  
  35.     // starting the timer here so that we include the cost of  
  36.     // all of the operations on the GPU.  if the data were  
  37.     // already on the GPU and we just timed the kernel  
  38.     // the timing would drop from 74 ms to 15 ms.  Very fast.  
  39.     cudaEvent_t     start, stop;  
  40.     HANDLE_ERROR( cudaEventCreate( &start ) );  
  41.     HANDLE_ERROR( cudaEventCreate( &stop ) );  
  42.     HANDLE_ERROR( cudaEventRecord( start, 0 ) );  
  43.   
  44.     // allocate memory on the GPU for the file's data  
  45.     unsigned char *dev_buffer;  
  46.     unsigned int *dev_histo;  
  47.     HANDLE_ERROR( cudaMalloc( (void**)&dev_buffer, SIZE ) );  
  48.     HANDLE_ERROR( cudaMemcpy( dev_buffer, buffer, SIZE,  
  49.                               cudaMemcpyHostToDevice ) );  
  50.   
  51.     HANDLE_ERROR( cudaMalloc( (void**)&dev_histo,  
  52.                               256 * sizeof( int ) ) );  
  53.     HANDLE_ERROR( cudaMemset( dev_histo, 0,  
  54.                               256 * sizeof( int ) ) );  
  55.   
  56.     // kernel launch - 2x the number of mps gave best timing  
  57.     cudaDeviceProp  prop;  
  58.     HANDLE_ERROR( cudaGetDeviceProperties( &prop, 0 ) );  
  59.     int blocks = prop.multiProcessorCount;  
  60.     histo_kernel<<<blocks*2,256>>>( dev_buffer,  
  61.                                     SIZE, dev_histo );  
  62.       
  63.     unsigned int    histo[256];  
  64.     HANDLE_ERROR( cudaMemcpy( histo, dev_histo,  
  65.                               256 * sizeof( int ),  
  66.                               cudaMemcpyDeviceToHost ) );  
  67.   
  68.     // get stop time, and display the timing results  
  69.     HANDLE_ERROR( cudaEventRecord( stop, 0 ) );  
  70.     HANDLE_ERROR( cudaEventSynchronize( stop ) );  
  71.     float   elapsedTime;  
  72.     HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime,  
  73.                                         start, stop ) );  
  74.     printf( "Time to generate:  %3.1f ms\n", elapsedTime );  
  75.   
  76.     long histoCount = 0;  
  77.     for (int i=0; i<256; i++) {  
  78.         histoCount += histo[i];  
  79.     }  
  80.     printf( "Histogram Sum:  %ld\n", histoCount );  
  81.   
  82.     // verify that we have the same counts via CPU  
  83.     for (int i=0; i<SIZE; i++)  
  84.         histo[buffer[i]]--;  
  85.     for (int i=0; i<256; i++) {  
  86.         if (histo[i] != 0)  
  87.             printf( "Failure at %d!\n", i );  
  88.     }  
  89.   
  90.     HANDLE_ERROR( cudaEventDestroy( start ) );  
  91.     HANDLE_ERROR( cudaEventDestroy( stop ) );  
  92.     cudaFree( dev_histo );  
  93.     cudaFree( dev_buffer );  
  94.     free( buffer );  
  95.     return 0;  
  96. }  


注：本文是作者对GPU高性能编程CUDA实战的学习总结。此书的代码可以在下面的链接下载，无需积分哦！

http://download.csdn.net/detail/celerychen2009/6360573


http://blog.csdn.net/celerychen2009/article/details/11318087
淘宝上买被子靠谱吗 淘宝上买被子靠谱吗
原创粉丝点击
热门IT博客
2016网络最红的歌曲2016网络最红的歌曲
资历章军龄算法资历章军龄算法
qq图标点亮软件
魔镜数据好用吗
易视网络电视在线直播
数据分析毕业论文题目
unicode转中文 java
如何编辑mac桌面图标
blued软件有时候卡
差评师数据设置
免费的出纳软件
linux认证
mac口红flamingo
最近网络流行词汇
英孚网络课程网址
迪拜旅游攻略大全网络
淘宝怎么改评价
星际淘宝网笔下
sql修改列默认值
学术猫数据库
热门问题 老师的惩罚 人脸识别 我在镇武司摸鱼那些年 重生之率土为王 我在大康的咸鱼生活 盘龙之生命进化 天生仙种 凡人之先天五行 春回大明朝 姑娘不必设防，我是瞎子 海马泡酒有什么功效 泡酒用什么酒好 桂花泡酒功效 桂花泡酒的功效 枸杞泡酒的功效 枸杞泡酒比例 灵芝泡酒有什么功效 猕猴桃泡酒功效与作用 石榴泡酒的做法 灵芝泡酒什么功效 蜈蚣泡酒功效与作用 板栗泡酒功效与作用 鲜人参怎么泡酒 240种泡酒方子 桃金娘泡酒有什么功效 三七泡酒的配方 海马枸杞泡酒配方 红枣枸杞泡酒配方 枸杞泡酒的比例 养生泡酒药材配方 黑枸杞泡酒配方 中药泡酒配方大全 枸杞泡酒配方 药材泡酒配方 海马泡酒配方 小老鼠泡酒有什么功效 枸杞红枣泡酒 泡酒料配方表 蝎子泡酒有什么功效 海马泡酒加什么药材 经典的水果泡酒配方 泡酒蛇20年咬人事件 蝎子泡酒的配方 雪莲花泡酒功效与作用 泡酒用什么酒 蛇泡酒的功效 什么泡酒好喝又有营养 黄芪泡酒配方 水果泡酒大全 枸杞泡酒的功效与作用 樱桃泡酒有什么功效

程序博客网，程序员的互联网技术博客家园。csdn论坛精品 msdn技术资料都在这里