Cuda束表决函数(warp vote)

原创 2015年07月08日 16:01:49

CUDA束表决函数

束表决函数:简单的理解就是在一个warp内进行表决


__all(int predicate):指的是predicate与0进行比较,如果当前线程所在的Wrap所有线程对应predicate不为0,则返回1。

__any(int predicate):指的是predicate与0进行比较,如果当前线程所在的Wrap有一个线程对应的predicate值不为0,则返回1。

__ballot(int predicate):指的是当前线程所在的Wrap中第N个线程对应的predicate值不为0,则将整数0的第N位进行置位。

 

//
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "device_functions.h"
#include <stdio.h>
 
__global__ void vote_all(int *a, int *b, int n)
{
    int tid = threadIdx.x;
    if (tid > n)
    {
       return;
    }
    int temp = a[tid];
    b[tid] = __all(temp >100);
}
 
__global__ void vote_any(int *a, int *b, int n)
{
    int tid = threadIdx.x;
    if (tid > n)
    {
       return;
    }
    int temp = a[tid];
    b[tid] = __any(temp >100);
}
 
__global__ void vote_ballot(int *a, int *b, int n)
{
    int tid = threadIdx.x;
    if (tid > n)
    {
       return;
    }
    int temp = a[tid];
    b[tid] = __ballot(temp >100);
}
 
int main()
{
    int *h_a, *h_b, *d_a, *d_b;
    int n = 256, m = 10;
    int nsize = n * sizeof(int);
    h_a = (int *)malloc(nsize);
    h_b = (int *)malloc(nsize);
    for (int i = 0; i < n; ++i)
    {
       h_a[i] = i;
    }
    memset(h_b, 0, nsize);
    cudaMalloc(&d_a, nsize);
    cudaMalloc(&d_b, nsize);
    cudaMemcpy(d_a, h_a, nsize, cudaMemcpyHostToDevice);
    cudaMemset(d_b, 0, nsize);
    vote_all<< <1, 256 >> >(d_a, d_b, n);
    cudaMemcpy(h_b, d_b, nsize, cudaMemcpyDeviceToHost);
    printf("vote_all():");
    for (int i = 0; i < n; ++i)
    {
       if (!(i % m))
       {
           printf("\n");
       }
       printf("%d", h_b[i]);
    }
    printf("\n");
    vote_any<<<1, 256 >> >(d_a, d_b, n);
    cudaMemcpy(h_b, d_b, nsize, cudaMemcpyDeviceToHost);
    printf("vote_any():");
    for (int i = 0; i < n; ++i)
    {
       if (!(i % m))
       {
           printf("\n");
       }
       printf("%d", h_b[i]);
    }
    printf("\n");
    vote_ballot<< <1, 256 >> >(d_a, d_b, n);
    cudaMemcpy(h_b, d_b, nsize, cudaMemcpyDeviceToHost);
    printf("vote_ballot():");
    for (int i = 0; i < n; ++i)
    {
       if (!(i % m))
       {
           printf("\n");
       }
       printf("%d", h_b[i]);
    }
    printf("\n");
}


#include <iostream>
 
using namespace std;
 
int main()
{
    int state = 0;
    int start = 10;
    for (int i = start; i <32; ++i)
    {
       state |= (1<< i);
    }
    cout<< state<< endl;
}


 

置位可以用或操作符“|”实现:y = x | (1 << n)  对x的第n位进行置位

清楚可以用与操作符”&“实现:y = x & (~(1 << n))

取反可以用异或操作符”^“实现: y = x ^ (1 << n)

Bit提取操作: bit = (x | (1 << n)) >> n;

版权声明:本文为博主原创文章,未经博主允许不得转载。

相关文章推荐

Use the CUDA Warp Watch

Use the CUDA Warp Watch VS2010的局部变量和全局变量工具窗口只支持在一个thread中查看变量一次,Nsight Debuger使用current focus...

CUDA中grid、block、thread、warp与SM、SP的关系

首先概括一下这几个概念。其中SM(Streaming Multiprocessor)和SP(streaming Processor)是硬件层次的,其中一个SM可以包含多个SP。thread是一个线程...

SimpleCV中shear()函数和warp()函数的区别

研究了SimpleCV中shear()函数和warp()函数的区别

CUDA C/C++关键字和函数高亮显示

CUDA C/C++关键字和函数高亮显示: 在上面HelloWorldCuda.cu文件中发现CUDA C/C++的关键字__global__等没有高亮显示,而且还有下划曲线。下面进行CUDA C/...

CUDA编程经验技术总结 系列之《核函数之线程技术》

0、核函数调用方式:kernel >>(); gridDim俗名线程格,是二维的(第三维为1),而blockDim是线程块,是三维的,cuda运行时容许启动一个二维线程格,并且线程格中的每个线程块都是...

CUDA 库函数-----cuRAND生成随机数

在主机侧使用随机数生成的库函数大致分为三个步骤: 1,声明随机数指针,指定随机数生成算法。 2,为生成算法设定初始值。 3,指定随机数生成个数,生成随机数,并保存在global me...

给cuda核函数传递二维数组的一种方法

#include /** * 需求:需要把若干个一维数组传给核函数 * 实现方法:在gpu生成一个一维的指针数组,每个元素指向一个普通一维数组。 * 把该指针数组的地址传递给核函数。 * ...

GPU编程自学4 —— CUDA核函数运行参数

深度学习的兴起,使得多线程以及GPU编程逐渐成为算法工程师无法规避的问题。这里主要记录自己的GPU自学历程。目录 《GPU编程自学1 —— 引言》 《GPU编程自学2 —— CUDA环境配置》 《GP...

研读OpenCV的GPU模块库函数,抠出自己CUDA代码之FAST角点检测篇

尝试利用CMake重编译OpenCV多次失败,苦于用不了OpenCV的GPU模块! 受一个师兄的启发,开始尝试从OpenCV库函数中抠出来自己的CUDA代码,忙活了一周终于有点起色。成功抠出来FAS...
内容举报
返回顶部
收藏助手
不良信息举报
您举报文章:深度学习:神经网络中的前向传播和反向传播算法推导
举报原因:
原因补充:

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