关闭

Thread与Warp

标签: warpcuda线程asm
512人阅读 评论(0) 收藏 举报
分类:

__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);
}




0
0

猜你在找
【直播】机器学习&数据挖掘7周实训--韦玮
【套餐】系统集成项目管理工程师顺利通关--徐朋
【直播】3小时掌握Docker最佳实战-徐西宁
【套餐】机器学习系列套餐(算法+实战)--唐宇迪
【直播】计算机视觉原理及实战--屈教授
【套餐】微信订阅号+服务号Java版 v2.0--翟东平
【直播】机器学习之矩阵--黄博士
【套餐】微信订阅号+服务号Java版 v2.0--翟东平
【直播】机器学习之凸优化--马博士
【套餐】Javascript 设计模式实战--曾亮
查看评论
* 以上用户言论只代表其个人观点,不代表CSDN网站的观点或立场
    个人资料
    • 访问:8040次
    • 积分:174
    • 等级:
    • 排名:千里之外
    • 原创:9篇
    • 转载:6篇
    • 译文:0篇
    • 评论:0条
    文章分类