Thread与Warp

原创 2015年07月10日 14:46:51

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

__popc(ballot(int predicate)):返回warp中bool不为零的线程数目

asm("mov.u32 %0, %laneid;" : "=r"(ret)):获得ret为当前线程在所在Warp中的ID

unsigned int ret;
asm("mov.u32 %0, %lanemask_lt;" : "=r"(ret));

__popc(ret & __ballot(int predicate)):返回的值为当前线程在所在的Warp中是第几个满足条件的


Example:

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "device_functions.h"
#include <iostream>

using namespace std;

__device__ __forceinline__ int laneId()
{
	unsigned int ret;
	asm("mov.u32 %0, %laneid;" : "=r"(ret));
	return ret;
}

__device__ __forceinline__ int laneMaskLt()
{
	unsigned int ret;
	asm("mov.u32 %0, %lanemask_lt;" : "=r"(ret));
	return ret;
}

__global__ void testKernel(int *a, int *b, int *c, int *d, int *e, int n)
{
	int x = threadIdx.x + blockIdx.x * blockDim.x;
	if (x >= n)
	{
		return;
	}
	a[x] = __ballot(x > 10);
	b[x] = laneMaskLt();
	d[x] = __popc(b[x] & a[x]);
	c[x] = __popc(a[x]);
	e[x] = laneId();
}

int main()
{
	int *a, *b, *c, *d, *e, *dev_a, *dev_b, *dev_c, *dev_d, *dev_e;
	int n = 64;
	int size = n * sizeof(int);
	a = (int *)malloc(size);
	b = (int *)malloc(size);
	c = (int *)malloc(size);
	d = (int *)malloc(size);
	e = (int *)malloc(size);
	cudaMalloc(&dev_a, size);
	cudaMalloc(&dev_b, size);
	cudaMalloc(&dev_c, size);
	cudaMalloc(&dev_d, size);
	cudaMalloc(&dev_e, size);

	testKernel<<<1, n>>>(dev_a, dev_b, dev_c, dev_d, dev_e, n);

	cudaMemcpy(a, dev_a, size, cudaMemcpyDeviceToHost);
	cudaMemcpy(b, dev_b, size, cudaMemcpyDeviceToHost);
	cudaMemcpy(c, dev_c, size, cudaMemcpyDeviceToHost);
	cudaMemcpy(d, dev_d, size, cudaMemcpyDeviceToHost);
	cudaMemcpy(e, dev_e, size, cudaMemcpyDeviceToHost);

	for (int i = 0; i < n; ++i)
	{
		printf("%d    %d    %d    %d    %d\n", a[i], b[i], c[i], d[i], e[i]);
	}
	cudaFree(dev_a);
	cudaFree(dev_b);
	cudaFree(dev_c);
	cudaFree(dev_d);
	cudaFree(dev_e);
	free(a);
	free(b);
	free(c);
	free(d);
	free(e);
}




Cuda 学习教程(五):GPU架构-Sp,sm,thread,block,grid,warp

SP(streaming Process),SM(streaming multiprocessor)是硬件(GPU hardware)概念。而thread,block,grid,warp是软件上的(C...

【并行计算-CUDA开发】CUDA编程——GPU架构,由sp,sm,thread,block,grid,warp说起

掌握部分硬件知识,有助于程序员编写更好的CUDA程序,提升CUDA程序性能,本文目的是理清sp,sm,thread,block,grid,warp之间的关系。由于作者能力有限,难免有疏漏,恳请读者批评...

CUDA编程系列--GPU架构,由sp,sm,thread,block,grid,warp说起

掌握部分硬件知识,有助于程序员编写更好的CUDA程序,提升CUDA程序性能,本文目的是理清sp,sm,thread,block,grid,warp之间的关系。由于作者能力有限,难免有疏漏,恳请读者批评...

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

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

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

首先概括一下这几个概念。其中SM(Streaming Multiprocessor)和SP(streaming Processor)是硬件层次的,其中一个SM可以包含多个SP。thread是一个线程,...
  • wvh2007
  • wvh2007
  • 2015年11月17日 18:30
  • 1831

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

[原]CUDA中grid、block、thread、warp与SM、SP的关系 2015-3-27阅读209 评论0 首先概括一下这几个概念。其中SM(Streaming Multiproce...

Warp : Haskell 的高性能 Web 服务器(译文)

Warp : Haskell 的高性能 Web 服务器(译文) 按 GHC 7.8 马上就要发布了。一个很大的改进就是加入了本文所说的并行 IO 管理器。从此之后 Haskell ...
  • fishmai
  • fishmai
  • 2016年05月30日 20:37
  • 1760

Cuda束表决函数(warp vote)

CUDA束表决函数 束表决函数:简单的理解就是在一个warp内进行表决 __all(int predicate):指的是predicate与0进行比较,如果当前线程所在的Wrap所有线程对应pred...

CUDA性能优化----warp深度解析

CUDA性能优化----warp深度解析   2017-01-12 16:41:07|  分类: HPC&CUDA优化 |  标签:gpu  cuda  hpc   |举报 |字号 订阅 ...

CUDA2.1-原理之索引与warp

本小节来自《大规模并行处理器编程实战》第四节,该书是很好的从内部原理结构上来讲述了CUDA的,对于理解CUDA很有帮助,借以博客的形式去繁取间,肯定会加入自己个人理解,所以有错误之处还望指正。 一、块...
内容举报
返回顶部
收藏助手
不良信息举报
您举报文章:Thread与Warp
举报原因:
原因补充:

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