extern __shared__

http://blog.csdn.net/openhero/article/details/3890578

首先是讲一下shared的memory的两种使用方法,然后讲解一下shared memory的bank conflict的问题,这个是shared memory访问能否高效的问题所在;

Shared memory的常规使用:

1. 使用固定大小的数组:

/************************************************************************/

/* Example */

/************************************************************************/

__global__ void shared_memory_1(float* result, int num, float* table_1)

{

__shared__ float sh_data[THREAD_SIZE];

int idx = threadIdx.x;

float ret = 0.0f;

sh_data[idx] = table_1[idx];

for (int i = 0; i < num; i++)

{

ret += sh_data[idx %BANK_CONFLICT];

}

result[idx] = ret;

}

这里的sh_data就是固定大小的数组;

2. 使用动态分配的数组:

extern __shared__ char array[];

__global__ void shared_memory_1(float* result, int num, float* table_1, int shared_size)

{

float* sh_data = (float*)array; // 这里就让sh_data指向了shared memory的第一个地址,就可以动态分配空间

float* sh_data2 = (float*)&sh_data[shared_size]; // 这里的shared_size的大小为sh_data的大小;

int idx = threadIdx.x;

float ret = 0.0f;

sh_data[idx] = table_1[idx];

for (int i = 0; i < num; i++)

{

ret += sh_data[idx %BANK_CONFLICT];

}

result[idx] = ret;

}

这里是动态分配的空间,extern __shared__ char array[];指定了shared的第一个变量的地址,这里其实是指向shared memory空间地址;后面的动态分配float* sh_data = (float*)array;让sh_data指向array其实就是指向shared memory上的第一个地址;

后面的float* sh_data2 = (float*)&sh_data[shared_size];这里的sh_data2是指向的第一个sh_data的shared_size的地址,就是sh_data就是有了shared_size的动态分配的空间;


3. 下面是讲解bank conflict

我们知道有每一个half-warp是16个thread,然后shared memory有16个bank,怎么分配这16个thread,分别到各自的bank去取shared memory,如果大家都到同一个bank取款,就会排队,这就造成了bank conflict,上面的代码可以用来验证一下bank conflict对代码性能造成的影响:

/************************************************************************/

/* Example */

/************************************************************************/

__global__ void shared_memory_1(float* result, int num, float* table_1)

{

__shared__ float sh_data[THREAD_SIZE];

int idx = threadIdx.x;

float ret = 0.0f;

sh_data[idx] = table_1[idx];

for (int i = 0; i < num; i++)

{

//并行的思想

ret += sh_data[idx %BANK_CONFLICT];

}

result[idx] = ret;

}

// 1,2,3,4,5,6,7.....16

#define BANK_CONFLICT 16

这里的BANK_CONFLICT 定义为从1到16的大小,可以自己修改,来看看bank conflict对性能的影响;当BANK_CONFLICT为2的时候,就会通用有8个thread同时访问同一个bank,因为idx%2的取值只有2个0和1,所以16个都会访问bank0和bank1,以此类推,就可以测试整个的性能;

当然我们还可以利用16bank conflict,大家都访问同一个bank的同一个数据的时候,就可以形成一个broadcast,那样就会把数据同时广播给16个thread,这样就可以合理利用shared memory的broadcast的机会。

下面贴出代码,最好自己测试一下;

/********************************************************************

* shared_memory_test.cu

* This is a example of the CUDA program.

* Author: zhao.kaiyong(at)gmail.com

* http://blog.csdn.NET/openhero

* http://www.comp.hkbu.edu.hk/~kyzhao/

*********************************************************************/

#include

#include

#include

#include

// 1,2,3,4,5,6,7.....16

#define BANK_CONFLICT 16

#define THREAD_SIZE 16

/************************************************************************/

/* static */

/************************************************************************/

__global__ void shared_memory_static(float* result, int num, float* table_1)

{

__shared__ float sh_data[THREAD_SIZE];

int idx = threadIdx.x;

float ret = 0.0f;

sh_data[idx] = table_1[idx];

for (int i = 0; i < num; i++)

{

ret += sh_data[idx%BANK_CONFLICT];

}

result[idx] = ret;

}

/************************************************************************/

/* dynamic */

/************************************************************************/

extern __shared__ char array[];

__global__ void shared_memory_dynamic(float* result, int num, float* table_1, int shared_size)

{

float* sh_data = (float*)array; // 这里就让sh_data指向了shared memory的第一个地址,就可以动态分配空间

float* sh_data2 = (float*)&sh_data[shared_size]; // 这里的shared_size的大小为sh_data的大小;

int idx = threadIdx.x;

float ret = 0.0f;

sh_data[idx] = table_1[idx];

for (int i = 0; i < num; i++)

{

ret += sh_data[idx%BANK_CONFLICT];

}

result[idx] = ret;

}

/************************************************************************/

/* Bank conflict */

/************************************************************************/

__global__ void shared_memory_bankconflict(float* result, int num, float* table_1)

{

__shared__ float sh_data[THREAD_SIZE];

int idx = threadIdx.x;

float ret = 0.0f;

sh_data[idx] = table_1[idx];

for (int i = 0; i < num; i++)

{

ret += sh_data[idx % BANK_CONFLICT];

}

result[idx] = ret;

}

/************************************************************************/

/* HelloCUDA */

/************************************************************************/

int main(int argc, char* argv[])

{

if ( cutCheckCmdLineFlag(argc, (const char**) argv, "device"))

{

cutilDeviceInit(argc, argv);

}else

{

int id = cutGetMaxGflopsDeviceId();

cudaSetDevice(id);

}

float *device_result = NULL;

float host_result[THREAD_SIZE] ={0};

CUDA_SAFE_CALL( cudaMalloc((void**) &device_result, sizeof(float) * THREAD_SIZE));

float *device_table_1 = NULL;

float host_table1[THREAD_SIZE] = {0};

for (int i = 0; i < THREAD_SIZE; i++ )

{

host_table1[i] = rand()%RAND_MAX;

}

CUDA_SAFE_CALL( cudaMalloc((void**) &device_table_1, sizeof(float) * THREAD_SIZE));

CUDA_SAFE_CALL( cudaMemcpy(device_table_1, host_table1, sizeof(float) * THREAD_SIZE, cudaMemcpyHostToDevice));

unsigned int timer = 0;

CUT_SAFE_CALL( cutCreateTimer( &timer));

CUT_SAFE_CALL( cutStartTimer( timer));

shared_memory_static<<<1, THREAD_SIZE>>>(device_result, 1000, device_table_1);

//shared_memory_dynamic<<<1, THREAD_SIZE>>>(device_result, 1000, device_table_1, 16);

//shared_memory_bankconflict<<<1, THREAD_SIZE>>>(device_result, 1000, device_table_1);

CUT_CHECK_ERROR("Kernel execution failed/n");

CUDA_SAFE_CALL( cudaMemcpy(host_result, device_result, sizeof(float) * THREAD_SIZE, cudaMemcpyDeviceToHost));

CUT_SAFE_CALL( cutStopTimer( timer));

printf("Processing time: %f (ms)/n", cutGetTimerValue( timer));

CUT_SAFE_CALL( cutDeleteTimer( timer));

for (int i = 0; i < THREAD_SIZE; i++)

{

printf("%f ", host_result[i]);

}

CUDA_SAFE_CALL( cudaFree(device_result));

CUDA_SAFE_CALL( cudaFree(device_table_1));

cutilExit(argc, argv);

}

这里只是一个简单的demo,大家可以测试一下。下一章节会将一些shared memory的更多的特性,更深入的讲解shared memory的一些隐藏的性质;

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处:http://www.mzph.cn/news/313169.shtml

如若内容造成侵权/违法违规/事实不符,请联系多彩编程网进行投诉反馈email:809451989@qq.com,一经查实,立即删除!

相关文章

SiteServer CMS 新版本 V6.13(2019年11月1日发布)

欢迎来到 SiteServer CMS V6.13 版本&#xff0c;经过两个月的连续迭代开发&#xff0c;V6.13版本新增了几项重要功能&#xff0c;我们希望你会喜欢&#xff0c;一些关键的亮点包括&#xff1a;。新增功能及BUG 修复经过两个月的连续迭代开发&#xff0c;V6.13 版本新增了部分功…

CUDA的global内存访问的问题

http://blog.csdn.net/OpenHero/article/details/3520578 关于CUDA的global内存访问的问题&#xff0c;怎么是访问的冲突&#xff0c;怎样才能更好的访问内存&#xff0c;达到更高的速度。下面先看几张图&#xff0c;这些图都是CUDA编程手册上的图&#xff0c;然后分别对这些…

C# 8 新特性 - 异步流 Asynchronous Streams

异步流 Asynchronous Streams例子 这是一个很简单的控制台程序。它有一个NumberFactory&#xff0c;它可以根据传递的参数来产生一串数字&#xff08;IEnumerable<int>&#xff09;。然后在这个程序中把每个数字都打印出来&#xff0c;同时在前边显示出当前的线程ID。 这…

__syncthreads()

http://www.cnblogs.com/dwdxdy/p/3215136.html __syncthreads()是cuda的内建函数&#xff0c;用于块内线程通信. __syncthreads() is you garden variety thread barrier. Any thread reaching the barrier waits until all of the other threads in that block also reach i…

互联网50周年!这有它的一张“出生证明”

2019 年 10 月 29 日是互联网的 50 周年&#xff0c;50 年前(1969 年 10 月 29 日)&#xff0c;加州大学洛杉矶分校的计算机将一个只有两个字母(LO)的数据包发送到斯坦福研究所的计算机上&#xff0c;这是互联网史上的第一个数据包&#xff0c;从此开启互联网时代的第一步。 当…

Eltwise_layer简介

http://www.voidcn.com/blog/thy_2014/article/p-6117416.html common_layer&#xff1a; ArgMaxLayer类&#xff1b; ConcatLayer类&#xff1a; EltwiseLayer类&#xff1b; FlattenLayer类&#xff1b; InnerProductLayer类&#xff1b; MVNLayer类&#xff1b; SilenceLaye…

PowerBI 秒级实时大屏展示方案 全面助力双十一

双十一来了&#xff0c;你准备好了吗&#xff1f;不管你是否准备完毕&#xff0c;我们带来了全网首发的 PowerBI 秒级实时大屏展示方案&#xff0c;你可以直接用来展示双十一的实时状况。我们一步步来说明这个套件模板教程。真实效果功能如下&#xff1a;全实时展示 双十一 当天…

caffe基本函数

http://blog.csdn.net/seven_first/article/details/47378697目录 主要函数 caffe_cpu_gemm 函数caffe_cpu_gemv 函数caffe_axpy 函数caffe_set 函数caffe_add_scalar 函数caffe_copy 函数caffe_scal 函数caffeine_cup_axpby 函数caffe_add caffe_sub caffe_mul caffe_div 函数…

优化 .net core 应用的 dockerfile

优化 .net core 应用的 dockerfileIntro在给 .net core 应用的写 dockerfile 的时候一直有个苦恼&#xff0c;就是如果有很多个项目&#xff0c;在 dockerfile 里写起来就会很繁琐&#xff0c;有很多项目文件要 copy&#xff0c;dockerfile 还不支持直接批量复制项目文件&#…

Ubuntu 上不了网

http://askubuntu.com/questions/441619/how-to-successfully-restart-a-network-without-reboot-over-ssh 指令&#xff1a; sudo service network-manager restart

C# 8 新特性 - 静态本地方法

从C# 8 开始&#xff0c;本地方法就可以是静态的了。 与其他的本地方法不同&#xff0c;静态的本地方法无法捕获任何本地状态量。 直接看例子&#xff1a; 这段代码里有两个本地方法&#xff0c;他们分别对实例的一个字段和方法里的一个本地变量进行了修改操作&#xff0c;也就…

Caffe使用其他人代码CUDA不兼容问题

只要将别人的上层代码使用即可&#xff0c;底层CUDA还是用自己的版本。

​.NET手撸2048小游戏

前言2048是一款益智小游戏&#xff0c;得益于其规则简单&#xff0c;又和 2的倍数有关&#xff0c;因此广为人知&#xff0c;特别是广受程序员的喜爱。本文将再次使用我自制的“准游戏引擎” FlysEngine&#xff0c;从空白窗口开始&#xff0c;演示如何“手撸” 2048小游戏&…

cannot import caffe

遇到问题“cannot import caffe”首先make pycaffe&#xff0c; 如果不行&#xff0c;分别尝试&#xff1a; 1. export PYHTONPATH/path/to/caffe/python:$PYTHONPATH 注&#xff1a;以上方法只能在当前terminal中生效&#xff0c;如需长久生效&#xff0c;需修改~/.bashrc …

自行实现高性能MVC

wcf虽然功能多、扩展性强但是也面临配置忒多&#xff0c;而且restful的功能相当怪异&#xff0c;并且目前没法移植。asp.net core虽然支持webapi&#xff0c;但是功能也相对繁多、配置复杂。就没有一个能让码农们安安心心的写webapi&#xff0c;无需考虑性能、配置、甚至根据问…

caffe matio问题

http://blog.csdn.net/houqiqi/article/details/46469981 注&#xff1a;如果指令行模式实在解决不了/lib/libcaffe.so: undefined reference to Mat_VarReadDataLinear问题&#xff0c;可以尝试在QT下进行训练和测试。 1&#xff0c; 下载matio(http://sourceforge.NET/pro…

技术管理者怎样跳出“泥潭”

近几年面试了不少新人&#xff0c;当问到职业规划时&#xff0c;大多都会说先积累技术&#xff0c;然后往架构师的方向发展。这可能是技术人的一个特质&#xff0c;喜欢跟机器相处&#xff0c;沉浸在代码之中&#xff0c;而不喜欢跟人打交道。现实的情况是&#xff0c;一些中小…

你或许以为你不需要领域驱动设计

作者&#xff1a;邹溪源&#xff0c;长沙资深互联网从业者&#xff0c;架构师社区合伙人&#xff01;一犹记得刚刚参加工作时&#xff0c;是地图厂商四维图新集团旗下的一家子公司&#xff0c;主要从事规划测绘相关软件研发的公司。当时我的项目是为勘测设计院提供相对应的应用…

undefined reference to Mat_VarCreate'

我们之前所写的方法&#xff0c;发现只能在QT编译模式下解决这个问题&#xff0c;而在指令行模式下不能编译通过&#xff0c;因为QT模式基于cmake&#xff0c;而make all基于Makefile 因而我们参照DeepLab的Makefile文件&#xff1a; 发现只要将其中关于matio的内容加到我们的…

redis为什么这么火该怎么用

最近一些人在介绍方案时&#xff0c;经常会出现redis这个词&#xff0c;于是很多小伙伴百度完redis也就觉得它是一个缓存&#xff0c;然后项目里面把数据丢进去完事&#xff0c;甚至有例如将实体属性拆分塞进redis hash里面的奇怪用法等等&#xff01;原因是什么呢&#xff1f;…