从几个简单例子了解CUDA内核的几个…

原创 2016年08月29日 12:50:06
首先来看向量相加从几个简单例子了解CUDA内核的几个重要的特性

下面是GPU 上向量相加的并行方法
1、GPU 上短长度的向量相加
从几个简单例子了解CUDA内核的几个重要的特性


再来看内积的kernel
__global__ void dot( int *a, int *b, int *c ) {
__shared__ int temp[THREADS_PER_BLOCK];     //每一个block内共享内存
index = threadIdx.x + blockIdx.x * blockDim.x;     //j计算线程的索引号。      
temp[threadIdx.x] = a[index] * b[index];  //每个线程分别计算乘积
__syncthreads();  //同步,等待所有乘积都计算好了
if( 0 == threadIdx.x ) {
  //由于每个block共享内存,所以在一个block 内执行一次求和。选0是因为每个block都至少有一个线程0吧
int sum = 0;
for( int i = 0; i < THREADS_PER_BLOCK; i++ )
sum += temp[i];
atomicAdd( c , sum );  //原子操作
}
}

main中 使用的方法
#define N (2048*2048)   //向量的总长度
#define THREADS_PER_BLOCK 512

dot<<< N+(THREADS_PER_BLOCK-1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK >>>( dev_a, dev_b, dev_c );

注意  <<>>

对于任意长度的向量内积。仿照任意长度向量相加,
http://blog.csdn.net/zhanglei0107/article/details/7328347

//计算向量的内积程序

#include
#define imin(a,b)    (a

//N为输入的向量的规模
const int N=33*1024;
const int threadsPerBlock=256;
const int blocksPerGrid=
imin(32,(N+threadsPerBlock-1)/threadsPerBlock);

__global__ void dot(float *a,float *b,float *c)
{
//每一个块上都有cache变量的拷贝,相互之间不影响
__shared__ float cache[threadsPerBlock];
//tid为线程的偏移量
int tid=threadIdx.x+blockIdx.x*blockDim.x;
int cacheIndex=threadIdx.x;
float temp=0;
while(tid
{
temp+=a[tid]*b[tid];
//增加的下标量为进程总数
tid+=blockDim.x*gridDim.x;

}
cache[cacheIndex]=temp;
//同步化当前块上的线程
__syncthreads();

int i=blockDim.x/2;
//在块内计算部分和
while(i!=0)
{
if(cacheIndex
cache[cacheIndex]+=cache[cacheIndex+i];
__syncthreads();
i/=2;
}

if(cacheIndex==0)
c[blockIdx.x]=cache[0];
}
橙色这段用了归约法




int main(void)
{

float *a,*b,c,*partial_c;

float *dev_a,*dev_b,*dev_partial_c;

a=(float*)malloc(N*sizeof(float));
b=(float*)malloc(N*sizeof(float));
partial_c=(float*)malloc(blocksPerGrid*sizeof(float));

cudaMalloc((void**)&dev_a,N*sizeof(float));
cudaMalloc((void**)&dev_b,N*sizeof(float));
cudaMalloc((void**)&dev_partial_c,blocksPerGrid*sizeof(float));


for(int i=0;i
{

a[i]=i;
b[i]=i*2;
}

cudaMemcpy(dev_a,a,N*sizeof(float),cudaMemcpyHostToDevice);
cudaMemcpy(dev_b,b,N*sizeof(float),cudaMemcpyHostToDevice);

dot<<>>(dev_a,dev_b,dev_partial_c);

cudaMemcpy(partial_c,dev_partial_c,blocksPerGrid*sizeof(float),cudaMemcpyDeviceToHost);

c=0;
for(int i=0;i
c+=partial_c[i];
    //此程序相当于计算0,1,...N-1的平方和
#define sum_squares(x)   (x*(x+1)*(2*x+1)/6)
    //测试向量内积的正确性
printf("Does GPU value %.6g = %.6g?\n",c,
2*sum_squares((float)(N-1)));


cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_partial_c);

free(a);
free(b);
free(partial_c);
//测试输出
int j;
scanf("%d",&j);






}

直方图的kernel
__global__ void histo_kernel( unsigned char *buffer,
                              long size,
                              unsigned int *histo ) {

    // clear out the accumulation buffer called temp
    // since we are launched with 256 threads, it is easy
    // to clear that memory with one write per thread
    __shared__  unsigned int temp[256];
    temp[threadIdx.x] = 0;  //初始化共享内存
    __syncthreads();

    // calculate the starting index and the offset to the next
    // block that each thread will be processing
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;
    while (i < size) {
        atomicAdd( &temp[buffer[i]], 1 );
        i += stride;
    }
   
__syncthreads();
atomicAdd(&(histo[threadIdx.x]), temp[threadIdx.x]);
    }


在GPU上多线程同时对数据集中的数据进行统计。为了保证计算过程中只有一个线程在对计算结果做改动,CUDA C使用原子操作的方式。







相关文章推荐

几个简单的jsp例子

  • 2013年10月20日 11:01
  • 153KB
  • 下载

几个简单的小例子学习c语言

  • 2013年05月23日 17:35
  • 793B
  • 下载

递归(一)几个简单的递归例子

刚接触递归的同学,可能难以理解递归,难以理解的点可能很多,例如: 1.函数为什么可以在自己的内部又调用自己呢? 2.既然可以自己调用自己,那么递归运行过程中一定回有很多层相互嵌套,到底什么时候不再嵌套...

c#几个简单的的例子

  • 2008年11月02日 12:45
  • 4.66MB
  • 下载

【Boost】boost库asio详解8——几个TCP的简单例子

摘于boost官网的几个例子, 做了点小修改, 笔记之. 同步客户端 void test_asio_synclient() { typedef boost::asio::io_service...
  • huang_xw
  • huang_xw
  • 2013年01月18日 22:56
  • 25223

几个简单的 ext初学例子

  • 2011年03月11日 15:17
  • 1MB
  • 下载

几个简单的博弈例子

几个简单的博弈例子 博弈的本意是下棋,现在也引申为:在一定条件下,遵守一定的规则,一个或几个拥有绝对理性思维的人或团队,从各自允许选择的行为或策略进行选择并加以实施,并从中各自取得相应结果或收益的过程...

C++中提供了多种基本的数据类型。实际上,这些远不能满足我们的需求,如复数(第10章的例子大多是处理虚数的),再如分数。本任务将设计一个简单的分数类,完成对分数的几个运算。一则巩固基于对象编程的方法,

#include using namespace std; class CFraction {private: int nume; // 分子 int deno; // 分...

flume-ng安装,应用,几个简单测试例子!

flume-ng安装,应用,几个简单测试例子!
内容举报
返回顶部
收藏助手
不良信息举报
您举报文章:从几个简单例子了解CUDA内核的几个…
举报原因:
原因补充:

(最多只允许输入30个字)