从几个简单例子了解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使用原子操作的方式。







CUDA学习,环境配置和简单例子

根据摩尔定律,每18个月,硬件的速度翻一番
  • helei001
  • helei001
  • 2014年04月07日 22:11
  • 1634

贪心算法-经典例子

package 二〇一七年三月二十四日; public class 贪心1 { public static void main(String[] args) { { int[] s...
  • qq_34789775
  • qq_34789775
  • 2017年06月08日 09:09
  • 349

GPU&CUDA几个基本概念

本文介绍 GPU 和 CUDA 相关的几个基本概念:SP、SM、warp、thread、block 和 grid。其中 SP,SM,warp 是硬件(GPU hardware)概念,而 thread,...
  • s_sunnyy
  • s_sunnyy
  • 2017年02月28日 17:09
  • 383

一个简单的CUDA程序

上一篇博客中讲到了如何下载并且搭建CUDA的开发环境,这次我将我第一次学到的CUDA程序记录下来,供自己和同行们日后查阅。 如果在安装CUDA之前安装了VisualStudio 20...
  • jiangcaiyang123
  • jiangcaiyang123
  • 2013年12月09日 23:53
  • 3600

CUDA例子

#include "cuda_runtime.h" #include "device_launch_parameters.h" #include #define N 10__global__ voi...
  • xuhang0910
  • xuhang0910
  • 2016年07月18日 13:38
  • 2564

第一章:一个简单的Web服务器

一些http基础知识,与怎样写一个自己的简单的web服务器
  • levycc
  • levycc
  • 2016年02月08日 21:25
  • 1386

几个简单php例子

1.is_函数 对变量进行种类判断 ()函数 强制转换 2.几个系统常量 其中line是魔术常量 运行结果 3.全局变量global 运行结果...
  • u012282037
  • u012282037
  • 2014年02月27日 10:53
  • 632

Java几个简单例子

例子1 Class.isPrimitive()方法:判断Class类是否是基础数据类型 Class.isAssignableFrom:用来判断一个类Class1和另一个类Class2是否相同或是另...
  • u013485533
  • u013485533
  • 2015年02月10日 16:59
  • 313

sql存储过程几个简单例子

导读:sql存储是数据库操作过程中比较重要的一个环节,对于一些初学者来说也是比较抽象难理解的,本文我将通过几个实例来解析数据库中的sql存储过程,这样就将抽象的事物形象化,比较容易理解。 例1:...
  • qq819632865
  • qq819632865
  • 2017年01月06日 01:23
  • 43

算法课笔记系列(一)—— 分治算法

首先必须强调的是,我是个算法渣,多少年了~~还是渣/(ㄒoㄒ)/~~ 这学期在上英文算法课,机缘巧合选了英文,觉得老师讲得还不错,所以没换到中文。想要通过这样总结一下加深算法理解,考试时也方便复习。 ...
  • Ying_Xu
  • Ying_Xu
  • 2016年03月16日 22:54
  • 2402
内容举报
返回顶部
收藏助手
不良信息举报
您举报文章:从几个简单例子了解CUDA内核的几个…
举报原因:
原因补充:

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