关于CUDA共享内存的使用

共享内存是定义在kernel函数里面还是外面?

我下面有一个类似图像直方图统计功能的内核函数,每个线程对应一个像素点。

如果不使用共享内存是这样的:

atomicAdd(&gpu_EO_0_stats_pix_count[catagory_row_id*LCU_total + LCU_id_ab], 1);
atomicAdd(&gpu_EO_0_stats_E[catagory_row_id*LCU_total + LCU_id_ab], e);

现在我希望使用共享内存,对共享内存原子写,在将最后结果汇聚到全局内存:

__shared__ stats_count shared_count[CATA_NUM];
__shared__ stats_E shared_E[CATA_NUM];

__syncthreads();

atomicAdd(&shared_count[catagory_row_id], 1);
atomicAdd(&shared_E[catagory_row_id], e);

__syncthreads();

            if (threadIdx.x == 0)//每个线程块的第一个线程
            {
                atomicAdd(&gpu_EO_0_stats_pix_count[catagory_row_id*LCU_total + LCU_id_ab], shared_count[catagory_row_id]);
                atomicAdd(&gpu_EO_0_stats_E[catagory_row_id*LCU_total + LCU_id_ab], shared_E[catagory_row_id]);
            }

但这样会出错,error code 77。

我想请问下出错的原因是什么?共享内存到底应该怎么定义和使用?

0

1个回答

0
lryanch
lryanch 呃。。。那篇文章好像没讲到什么关键信息吧
3 年多之前 回复
Csdn user default icon
上传中...
上传图片
插入图片
抄袭、复制答案,以达到刷声望分或其他目的的行为,在CSDN问答是严格禁止的,一经发现立刻封号。是时候展现真正的技术了!
其他相关推荐
CUDA学习之路2 共享内存的使用
这两天在写一个对任意长度数组进行扫描的cuda代码,由于对共享内存的理解不清楚,浪费了好多时间,写篇博客记录一下,避免下次再犯。 正文: 当我们使用共享内存时,要谨记,但凡涉及到共享内存的操作,一定要记得使用线程同步__syncthread();否则必然会造成数据的计算出错。因为共享内存在块内对所有线程都是可见的,是大家一起使用的。如果你在每次使用之后不同步一下,那么当运算的快的线程运算完成后...
CUDA学习--内存处理之共享内存(3)
1.共享内存共享内存实际上是可受用户控制的一级缓存。每个SM中的一级缓存与共享内存共享一个64KB的内存段。在实际中,共享内存的速度几乎在所有的GPU中都一致(大约为1.5TB/s的带宽),因为共享内存的速度受核时钟频率驱动。因此在任何显卡中,无论是否为高端显卡,除了使用寄存器外,还要更有效的使用共享内存。然而,GPU执行的是一种内存的加载-存储模型,即所有的操作都要在指令载入寄存器后才能执行。因此
CUDA学习笔记(十二)共享内存
CUDA SHARED MEMORY shared memory在之前的博文有些介绍,这部分会专门讲解其内容。在global Memory部分,数据对齐和连续是很重要的话题,当使用L1的时候,对齐问题可以忽略,但是非连续的获取内存依然会降低性能。依赖于算法本质,某些情况下,非连续访问是不可避免的。使用shared memory是另一种提高性能的方式。 GPU上的memory有两种: · On...
CUDA优化实例(三)共享内存
CUDA优化实例(三)共享内存 前言 经过前面的实验发现,共享内存是优化CUDA程序的核心方法。共享内存可以通过对全局内存数据进行合并访问,让kernel内交错的内存需求去访问共享内存。如:矩阵转置问题,将二维内存的行写入二维内存的列。对列的写入就是一个内存交错访问的例子。可以用合并的方式将块要操作的数据写入共享内存,让复杂的内存交错访问访问共享内存,然后将结果以合并的方式写入全局内存。 ...
CUDA学习笔记(6) 共享内存与全局内存
共享内存(Shared memory)是位于每个流处理器组(SM)中的高速内存空间,主要作用是存放一个线程块(Block)中所有线程都会频繁访问的数据。流处理器(SP)访问它的速度仅比寄存器(Register)慢,它的速度远比全局显存快。但是他也是相当宝贵的资源,一般只有几十KByte,  这里以我电脑上的“Quadro K620”为例: 硬件资源 参数 流处理器(SP) 128 *
CUDA编程(七)共享内存与Thread的同步
CUDA编程(七)共享内存与Thread的同步在之前我们通过block,继续增大了线程的数量,结果还是比较令人满意的,但是也产生了一个新的问题,即,我们在CPU端的加和压力变得很大,所以我们想到能不能从GPU上直接完成这个工作。我们知道每个block内部的Thread之间是可以同步和通讯的,本篇我们将让每个block把每个thread的计算结果进行加和。所以本篇博客我们将研究CUDA架构中Threa
CUDA矩阵乘法——利用共享存储器
上篇的方法是在全局存储区中,这样对取数据时速度回很慢,影响性能,而设备中线程对块中的共享存储区中数据读取时速度是很快的,并且在全局存储区中进行读取时,有很多数组元素的重复读取。因此,先将需要计算的数组数据读取到共享存储区中,再利用共享存储区中的数据进行计算,就会提高性能。但由于每个块的共享存储区的存储空间一般很小,以本人8400MG为例,只有16KB,因此在一个块内需要的数据量大时,有必要对数据进
CUDA 共享内存的动态分配
__global__ static void sumOfSquares(int *num,int *result,clock_t *time) { extern __shared__ int sharedA[]; extern __shared__ int sharedB[];    } 这个就是动态分配 其中每个shared分配的
CUDA共享内存操作(__shared__关键字)
CUDA 对GPU的显存操作时,若线程频繁对某个数据进行读写操作,可以设置操作的数据常驻缓存,这样可以进一步提高代码的运行效率,并且同一个线程块内的所有线程共享该内存区域。然而当出现多个线程对同一个内存区域进行操作时,需要对线程进行同步操作,从而避免竞争的发生。cuda中使用__shared__关键字,这里使用__syncthreads()控制线程同步。 如核函数dot所示,代码生成一个数组并计算...
CUDA总结:共享内存
共享内存是片上内存(on-chip),所以速度比一般的显存快很多,如(全局内存、常亮内存、纹理内存)。共享内存是gpu中,带宽仅次于寄存器的存储器。共享内存是有限的,与L1 Cache公用一块on-chip内存,用户可以调整L1 cache与共享内存的大小组合。在on-chip内存的基础上,共享内存还实现了“并行访存”:共享内存被划分为大小相等的n个部分(每个部分称为一个bank),同一时刻的n个访
CUDA之静态、动态共享内存分配详解
静态分配 加上前缀 shared __shared__ int _ss[1024];1 动态分配 当我们在编程时,不清楚shared memory 数组开多大,就要用到动态分配。  分为两部分:  1, 声明 extern __shared__ int _s[];1 2, 在调用kernel 时加上数组的大小。 xxx_kernelgrid, block, sha
cuda共享内存,全局内存,纹理等的解释
开始阅读粗大资料感觉文献1的描述讲得明白,摘录到这里方便他人了解。 增加简单排版后,摘录 1.共享内存 目前 CUDA 装置中,每个 multiprocessor 有 16KB 的 shared memory。 Shared memory 分成 16 个 bank。如果同时每个 thread 是存取不同的 bank,就不会产生任何问题,存取 shared memory 的速度和存取寄
cuda实现二维卷积共享内存
本次cuda课作业为实现cuda卷积。核心代码为 1.cpu卷积代码 void Conv2(float** filter, float** arr, float** res, int filter_size, int arr_size) { int temp; for (int i = 0; i<arr_size; i++) { for (int j = 0; j<ar...
cuda shared(共享内存)—一个容易错误的点
可能出现的代码情况: template __global__ void kCopy(int row_size, float **col_valuess, float *row_values) { int row_index = blockIdx.x * blockDim.x + threadIdx.x; if (row_index >= row_size) return; vol
CUDA矩阵转置(共享内存 tile)
Udacity的CUDA编程课程中介绍了CUDA实现矩阵转置的六种方式,本文介绍其中的一种方式 如果矩阵为N*N的方阵。该方式让每个线程处理一个矩阵元素,总共需要N*N个线程。首先,声明两个常量并配置blocks,threads: const int N=1024; const int K=32; dim3 blocks(N/K,N/K); dim3 threads(K,K); 内核函...
CUDA申请动态共享内存
CUDA申请动态共享内存 直接上代码: #include <stdio.h> #include <cuda_runtime.h> __global__ void kernel(){ int index = blockIdx.x*blockDim.x + threadIdx.x; extern __shared__ float s_int[]; ...
[菜鸟每天来段CUDA_C]基于共享内存的位图与syncthreads的使用
本文使用CUDA实现基于共享内存的位图显示。位图中每个位置的像素值由每个线程计算,计算结果保存到缓冲区(共享内存)中。 结果为一个由多个绿色球形构成的网格。
cuda编程:关于共享内存(shared memory)和存储体(bank)的事实和疑惑
主要是在研究访问共享内存会产生bank conflict时,自己产生的疑惑。对于这
【并行计算-CUDA开发】关于共享内存(shared memory)和存储体(bank)的事实和疑惑
关于共享内存(shared memory)和存储体(bank)的事实和疑惑 主要是在研究访问共享内存会产生bank conflict时,自己产生的疑惑。对于这点疑惑,网上都没有相关描述, 不管是国内还是国外的网上资料。貌似大家都是当作一个事实,一个公理,而没有对其仔细研究。还是我自己才学疏浅,不知道某些知识。 比如下面这篇讲解bank conflict的文章。
数组求和的快速方法(利用cuda的共享内存)--第一部分之源码分析
代码来自于这里 https://code.google.com/p/stanford-cs193g-sp2010/source/browse/trunk/tutorials/sum_reduction.cu
CUDA编程—通过shared memory优化矩阵相乘
使用shared memory带来两点优化: (1)使用共享存储器减少了线程块中的线程必须访问的数据总量。 (2)使用共享存储器来实现存储空间的合并。(下篇文章展开介绍)
CUDA内存(二) 共享内存 shared memory
其他: CUDA内存(一) 寄存器 目录共享内存使用共享内存排序:测试结果 共享内存 共享内存实际上是可受用户控制的一级缓存. [^1] 只有当数据重复利用, 全局内存合并, 或者线程之间有共享数据时, 使用共享内存才合适. 使用共享内存排序: SortArray.h #pragma once #include "Global.h" #include "device_launch_paramet...
CUDA 学习(十七)、优化策略2:内存因素
一、内存带宽        内存带宽和延迟是所有应用程序都要考虑的关键因素,尤其是GPU应用程序。带宽是指与某个给定目标之间传输的数据量。在GPU的情况下,我们主要关心的是全局内存带宽。延迟则是操作完成所用的时间。        GPU上的内存延迟设计为由运行在其他线程束中的线程所隐藏。当线程束访问的内存位置不可用时,硬件向内存提交一次读或者写请求。如果同一线程束上其他线程访问的是相邻内存位置
CUDA实现基于共享内存的位图显示
CUDA实现基于共享内存的位图显示,线程同步
数组求和的快速方法(利用cuda的共享内存)--第三部分之性能分析
测试的是
CUDA之矩阵转置程序优化实例
Catalog 已经达到极限了?影响代码性能的两个主要方面优化代码内存操作 看代码内存操作是否有效——DRAM utilizationcoalesce合并从little’s Law中找继续优化的方法SM中的occupancy占用率 优化代码计算性能 减小线程发散度选择效率更高的数学计算 已经达到极限了? 经过了对行、各元素并行化处理,OK,似乎目前得到了一个还不
CUDA的内存结构,通过实例展示寄存器和共享内存的使用
本章将介绍CUDA的内存结构,通过实例展示寄存器和共享内存的使用。CUDA内存结构GPU的内存结构和CPU类似,但也存在一些区别,GPU的内存中可读写的有:寄存器(registers)、Local memory、共享内存(shared memory)和全局内存(global memory),只读的有:常量内存(constant memory)和纹理内存(texture memory)。CUDA To
CUDA教程之——共享存储器(1)-矩阵相乘
利用矩阵乘法说明如何使用共享存储器   参考资料:1)CUDA_C_Programming_Guide;        2)CUDA_C_Best_Practice
CUDA矩阵乘法
一、使用全局内存 矩阵乘法,即用矩阵A每行与矩阵B的每列,依次做乘积累加就可以得到各个元素的值。在CPU上用三层循环实现。这里是将二维数组用一维的形式表示,即按行存储。 size_t size = WIDTH*WIDTH * sizeof(int); int *h_A = (int *)malloc(size); int *h_B = (int*)malloc(size); int *
GPU共享内存:pycuda使用教程
最近在学习pycuda,因为资料比较少,而且杂乱,所以我就拷贝到了自己的博客 使用Python写CUDA程序有两种方式: * Numba * PyCUDA numbapro现在已经不推荐使用了,功能被拆分并分别被集成到accelerate和Numba了。 例子 numba Numba通过及时编译机制(JIT)优化Python代码,Numba可以针对本机的硬件环境进行优化,同时支持CP...
CUDA学习--矩阵乘法的并行运算
1. CUDA学习步骤 CPU实现 a*b = c 的矩阵乘法(矩阵尺寸是n*m的,n和m大于1000) 下载 https://developer.nvidia.com/cuda-downloads,安装好cuda 将cpu代码移植到cuda。将CPU值传入GPU,使用cuda计算,与cpu结果对比。 优化思路1:将矩阵分块进行计算 优化思路2:使用share memory进行优化 优化思路3:将数
CUDA(Ⅴ):共享内存(shared memory)、同步(_syncthreads)、点积运算
截止现在:将Block分解为Thread的目的只是为了解决Block数量的硬件限制。 问题:CUDA C的关键字__share__添加到变量声明中,这将使这个变量驻留在shared memory中,这样做的目的是什么? CUDA C编辑器对shared memory中的变量与普通变量将分别采取不同的处理方式。对于GPU上启动的每个Block,CUDA C编译器都将创建该变量的一个副本。Bloc...
cuda中线程块共享存储(shared memory)加速较全局存储(global memory)之优势
(仅供参考,各方面表述可能有错,概不负责)           首先个人观点说明两个概念,(1)共享存储加速,就是一个线程块内不仅所有线程并发执行,而且各线程还通过共享的内存来实现协作,进一步提升加速效果;(2)全局存储加速,仅仅是前者的前一个加速功能,即所有线程并发执行,线程之间不存在相互协作关系。           考虑两个等大小正方矩阵A(m x m)与B(m x m)相乘,结果存储在
CUDA计算直方图(二) 共享内存 __shared__
共享内存 共享内存是一块特殊的内存, 因为它存在于芯片上并且存取速度比全局内存快. 可以在共享内存上创建一个包含256个bin的局部统计直方图, 最后将所有共享内存上计算得到的统计直方图通过原子操作汇总到全局内存. 这样可以节省存储直方图结果的时间. 下图是GTX1050 的内存容量. 分治 // 共享内存. __shared__ Cuda32u d_bin_data_shared[256]; ...
共享存储器bank conflict
对每个bank在每个周期只能进行一次32位的读写操作,因此也说每个bank的带宽为每周期32bit。   只要同一个 warp 的不同线程会访问到同一个 bank 的不同地址就会发生 bank conflict,除此之外的都不会发生 bank conflict。
CUDA使用笔记(一)矩阵乘法
简介: 本文介绍cublasSgemm()函数的使用。在c/c++中,通常我们将2维矩阵按行存储为一维数组。但是在显存中,矩阵是按列存储的。因此,我们在实际使用时,对cublasSgemm()中的各个参数的赋值可能会搞不清楚。 本文,以一个具体的矩阵乘法案例为例子,介绍cublasSgemm()函数的使用。     正文: 我们以下图所示的矩阵运算为例进行讲解。     因为...
cuda shared memory 静态分配和动态分配
静态分配加上前缀 shared__shared__ int _ss[1024];动态分配当我们在编程时,不清楚shared memory 数组开多大,就要用到动态分配。 分为两部分: 1, 声明extern __shared__ int _s[];2, 在调用kernel 时加上数组的大小。xxx_kernel<<<grid, block, sharedMemSize>>>();内存分布下面通过
CUDA教程之——共享存储器(2)-矩阵与自身转置相乘
用共享存储器存储子矩阵;首先确定线程块,然后交叉使用线程块里的行列索引
CUDA线程协作之共享存储器“__shared__”&amp;&amp;“__syncthreads()”
在GPU并行编程中,一般情况下,各个处理器都需要了解其他处理器的执行状态,在各个并行副本之间进行通信和协作,这涉及到不同线程间的通信机制和并行执行线程的同步机制。共享内存“__share__”CUDA中的线程协作主要是通过共享内存实现的。使用关键字“__share__”声明共享变量,将使这个变量驻留在共享内存中,该变量具有以下特征:位于线程块的共享存储器空间中与线程块具有相同的生命周期仅可通过块内...
共享内存: Shared Memory
Linux进程间通信方法有: 同一主机的父子进程间:管道 同一主机上的进程间:FIFO(命名管道)、消息队列、信号量、信号、共享内存,socket。 不同主机间:socket 这里总结共享内存知识,主要是代码总结。详细帮助见unp vol2 : 进程间通信,或者man 文档。System V: shmget shmat shmdt 示例代码:#include <stdio.h> #include
文章热词 机器学习教程 Objective-C培训 交互设计视频教程 颜色模型 设计制作学习
相关热词 mysql关联查询两次本表 native底部 react extjs glyph 图标 python深度学习cuda 关于大数据培训