#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <iostream>
#include <cuda.h>
#include"device_functions.h"
/*
线程块Block的数量限制是:65535
我的机器上maxThreadPerBlock = 1024.
也就是一个Block中最多1024个线程
---
本篇主要是共享变量的使用
*/
const int N = (33 * 1024);
#define imin(a,b) ( a < b ? a:b )
const int threadPerBlock = 512;
const int blockPerGrid = imin(32 , (N + threadPerBlock -1 ) / threadPerBlock );
__global__ void dot(float *a , float *b , float* c)
{
__shared__ float cache[threadPerBlock];
int tid = threadIdx.x + blockIdx.x * blockDim.x;//获得线程索引
// printf("tid = %d\n",tid);
int cacheIndex = threadIdx.x;
float temp = 0;
while( tid < N) //必须检查是否在正确的偏移范围类内
{
temp += a[ tid ] * b[ tid];
tid += blockDim.x * gridDim.x;
}
//设置cache中相应位置上的值
cache[ cacheIndex ] = temp;
//对线程块中的线程进行同步
__syncthreads();//每个线程都必须保证前面的操作执行完了之后才能执行下面的语句,
//也就是说,如果有的线程提前完成了前面的操作,而其他线程还没完成,那么该线程将进入挂起状态
//这其实是一种约定,大家都必须先把前面的X干完,才能一起进入下个阶段
//必须要这样保证,因为如果M提前进入下一个阶段的话,他将去取其他cache位置上的数,但是他会发现,有些数字根本还没装进来,因为其他线程还没计算完,如果他执意要进行工作,那么必定会导致最终的计算错误。所以,这里必须要同步,要大伙儿一起进入下一个阶段才能保证结果的正确。
//归约运算
int i = blockDim.x / 2;
while( i != 0)
{
if( cacheIndex < i )
cache[ cacheIndex ] += cache[ cacheIndex + i ];
__syncthreads();//和上面的解释相同道理,因为都是读cache数据,必须要保证所有线程的写操作都完成,才能继续下一轮的读取,再写,再读...
i /= 2;
}
if( cacheIndex == 0)
c[ blockIdx.x] = cache[0];
}
int main()
{
float *a , *b , c , *partial_c;
float *dev_a, *dev_b , *dev_partial_c;
//在CPU上分配内存
a = (float*)malloc( N * sizeof(float));
b = (float*)malloc( N * sizeof(float));
partial_c = (float*)malloc( blockPerGrid * sizeof(float));
//在GPU上分配内存
cudaMalloc( (void**)&dev_a , N * sizeof(float ) );
cudaMalloc( (void**)&dev_b , N * sizeof(float ) );
cudaMalloc( (void**)&dev_partial_c , blockPerGrid * sizeof(float ) );
//在CPU上为数组a,b赋初值
int i = 0;
for( i = 0; i < N; i++)
{
*(a+i) = i ;
*(b+i) = i ;
//a[i] = i;
//b[i] = i;
//printf("%f\n",*(a+i));
}
//将数组a, b 复制到GPU上去计算
cudaMemcpy( dev_a , a , N * sizeof( float ) , cudaMemcpyHostToDevice);
cudaMemcpy( dev_b , b , N * sizeof( float ) , cudaMemcpyHostToDevice);
dot<<<blockPerGrid,threadPerBlock>>>( dev_a , dev_b , dev_partial_c );//那么将会有128 * 128 个核函数运行
cudaMemcpy( partial_c , dev_partial_c , blockPerGrid * sizeof( float ) , cudaMemcpyDeviceToHost);
//在CPU上完成最终的求和运算
c = 0;
for(int i = 0 ; i < blockPerGrid ; i++)
{
c += *(partial_c + i );
}
cudaFree( dev_a );
cudaFree( dev_b );
cudaFree( dev_partial_c );
//释放CPU上的内存
free( a );
free( b );
free( partial_c );
printf("c = %f" , c);
getchar();
return 0;
}
CUDA-CODE6-共享和同步
最新推荐文章于 2024-07-31 03:06:45 发布