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;

cuda版本的word2vec

上篇博客的快排中用到了基于warp的cuda操作用于分隔数组, 为什么要将控制线程的级别定义为warp呢? 在一个warp内,线程的可以通过__ballot函数,并发的获取这32个数中于pivot的...
  • sunmenggmail
  • sunmenggmail
  • 2015-01-20 17:22:23
  • 2028

cuda编程实践

昨天晚上弄到12. 终于迈出了c++到cuda的第一步,经过测试,基本可以确定是真的实现了。表示很兴奋,可是别人并不懂,没地方去庆祝啊。。。。    废话不多说,下面来介绍一下基本的实现过程。(国人在...
  • u012839187
  • u012839187
  • 2015-06-07 16:35:35
  • 3062

Thread与Warp

__ballot(int predicate):指的是当前线程所在的Wrap中第N个线程对应的predicate值不为零,则将整数零的第N位进行置位 __popc(ballot(int predica...
  • u010646276
  • u010646276
  • 2015-07-10 14:46:51
  • 694

【并行计算-CUDA开发】CUDA ---- Warp解析

Warp 逻辑上,所有thread是并行的,但是,从硬件的角度来说,实际上并不是所有的thread能够在同一时刻执行,接下来我们将解释有关warp的一些本质。 Warps and Thread...
  • LG1259156776
  • LG1259156776
  • 2016-10-13 22:42:08
  • 587

Cuda学习笔记(一)——sm流处理器簇对blocks的调度策略

由于GPU目前在各行各业的广泛应用,无论是深度学习、大数据、云计算等都离不开GPU的并行加速,前阵子自学了Cuda-c编程,希望将来的研究工作能够用得上。   Cuda系列总共有4篇,这里主要用于记...
  • GH234505
  • GH234505
  • 2016-04-10 22:45:57
  • 3526

cuda __shfl_xor

__shfl是cuda提供的wrap(线程束)级别的方法,一般线程束为32。相似的方法有几个。 __shfl() Direct copy from indexed lane __shfl_up(...
  • qq_16097611
  • qq_16097611
  • 2016-06-04 17:52:03
  • 1934

vote

#include #include #include using namespace std; map candidaes; typedef map::iterator VstrItor; int...
  • libin88211
  • libin88211
  • 2015-07-17 22:51:00
  • 347

多数投票算法 --- A linear time majority vote algorithm

多数投票算法 --- A linear time majority vote algorithm
  • u012965373
  • u012965373
  • 2016-08-01 10:49:35
  • 844

束表决函数作用

warp vote是2.0+的一个特性,用来在warp内交换1-bit的信息的。 有三种, 一种是直接将32个1-bit的位,映射成1个32位整数。warp中的线程0,将占据这个映射结果的最低0位,线...
  • niefeng111
  • niefeng111
  • 2014-03-13 17:53:04
  • 671

Cuda束表决函数(warp vote)

CUDA束表决函数 束表决函数:简单的理解就是在一个warp内进行表决 __all(int predicate):指的是predicate与0进行比较,如果当前线程所在的Wrap所有线程对应pred...
  • u010646276
  • u010646276
  • 2015-07-08 16:01:49
  • 885
收藏助手
不良信息举报
您举报文章:Cuda束表决函数(warp vote)
举报原因:
原因补充:

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