- html - 出于某种原因,IE8 对我的 Sass 文件中继承的 html5 CSS 不友好?
- JMeter 在响应断言中使用 span 标签的问题
- html - 在 :hover and :active? 上具有不同效果的 CSS 动画
- html - 相对于居中的 html 内容固定的 CSS 重复背景?
我有 5 个大型数组 A(N*5)、B(N*5)、C(N*5)、D(N*5)、E(N*2)数字 5 和 2 表示这些变量在不同平面/轴中的分量。这就是我以这种方式构建数组的原因,这样我就可以在编写代码时可视化数据。N ~ 200^3 ~ 8e06 个节点
例如:这是我的内核最简单的形式,我在其中对全局内存进行所有计算。
#define N 200*200*200
__global__ void kernel(doube *A, double *B, double *C,
double *D, double *E, double *res1, double *res2,
double *res3, double *res4 )
{
int a, idx=threadIdx.x + blockIdx.x * blockDim.x;
if(idx>=N) {return;}
res1[idx]=0.; res2[idx]=0.;
res3[idx]=0.; res4[idx]=0.
for (a=0; a<5; a++)
{
res1[idx] += A[idx*5+a]*B[idx*5+a]+C[idx*5+a] ;
res2[idx] += D[idx*5+a]*C[idx*5+a]+E[idx*2+0] ;
res3[idx] += E[idx*2+0]*D[idx*5+a]-C[idx*5+a] ;
res4[idx] += C[idx*5+a]*E[idx*2+1]-D[idx*5+a] ;
}
}
我知道可以去掉“for”循环,但我把它留在这里,因为这样看代码很方便。这行得通,但显然即使在删除“for”循环后,它对于 Tesla K40 卡来说也是极其低效和缓慢的。 “for”循环中显示的算法只是为了提供一个想法,实际的计算要长得多,并且与 res1、res2... 混杂在一起。
我已经实现了以下改进有限,但是我想通过共享内存的过载进一步改进它。
#define THREADS_PER_BLOCK 256
__global__ void kernel_shared(doube *A, double *B, double *C,
double *D, double *E, double *res1, double *res2,
double *res3, double *res4 )
{
int a, idx=threadIdx.x + blockIdx.x * blockDim.x;
int ix = threadIdx.x;
__shared__ double A_sh[5*THREADS_PER_BLOCK];
__shared__ double B_sh[5*THREADS_PER_BLOCK];
__shared__ double C_sh[5*THREADS_PER_BLOCK];
__shared__ double D_sh[5*THREADS_PER_BLOCK];
__shared__ double E_sh[2*THREADS_PER_BLOCK];
//Ofcourse this will not work for all arrays in shared memory;
so I am allowed to put any 2 or 3 variables (As & Bs) of
my choice in shared and leave rest in the global memory.
for(int a=0; a<5; a++)
{
A_sh[ix*5 + a] = A[idx*5 + a] ;
B_sh[ix*5 + a] = B[idx*5 + a] ;
}
__syncthreads();
if(idx>=N) {return;}
res1[idx]=0.; res2[idx]=0.;
res3[idx]=0.; res4[idx]=0.
for (a=0; a<5; a++)
{
res1[idx] += A_sh[ix*5+a]*B_sh[ix*5+a]+C[idx*5+a];
res2[idx] += B_sh[ix*5+a]*C[idx*5+a]+E[idx*2+0] ;
res3[idx] += E[idx*2+0]*D[idx*5+a]-C[idx*5+a] ;
res4[idx] += B_sh[ix*5+a]*E[idx*2+1]-D[idx*5+a] ;
}
}
这有点帮助,但我想实现其中一项减少方法(没有银行冲突)来提高性能,我可以把所有我共享的变量(可能是平铺方法)然后进行计算部分。我在 CUDA_Sample 文件夹中看到了缩减示例,但是那个示例仅对共享中的一个 vector 求和,而不涉及共享内存中多个数组的任何复杂算术。我将不胜感激任何帮助或建议来改进我现有的 kernel_shared 方法以包括减少方法。
最佳答案
检查您的初始内核,我们注意到对于 a
的每个值,您在计算要加起来的四个增量时最多使用 12 个值(可能少于 12 个,我没有准确计数)。这一切都非常适合您的寄存器文件 - 即使是 double 值: 12 * sizeof(double) ,加上 4 * sizeof(double) 中间结果使每个线程有 32 个 4 字节寄存器。即使每个 block 有 1024 个线程,也远远超出了限制。
现在,你的内核运行缓慢的原因主要是
这是您可以在任何 CUDA 编程演示中读到的内容;我只是简单地说,不是每个线程自己处理几个连续的数组元素,而是应该将其交错在 warp 的 channel 之间,或者更好的是在 block 的线程之间。因此,而不是线程全局索引 idx 处理
5 * idx
5 * idx + 1
...
5 * idx + 4
让它处理
5 * blockDim.x * blockIdx.x + threadIdx.x
5 * blockDim.x * blockIdx.x + threadIdx.x + blockDim.x
...
5 * blockDim.x * blockIdx.x + threadIdx.x + 4 * blockDim.x
这样,每当线程读取或写入时,它们的读取和写入合并。在您的情况下,这可能有点棘手,因为您的某些访问模式略有不同,但您明白了。
这个问题更具体到你的情况。你看,你真的不需要在 每一个 添加后更改全局的 resN[idx]
值,而且你当然不关心阅读每当你要写的时候,它就在那里。正如您的内核所代表的那样,单个线程为 resN[idx]
计算一个新值 - 因此它可以将寄存器中的内容相加,然后写入 resN[idx]
当它完成时(甚至不看它的地址)。
如果您按照我在第 1 点中的建议更改内存访问模式,则实现第 2 点中的建议会变得更加棘手,因为您需要将同一 warp 中的多个 channel 的值相加,并且可能使确保您不会跨越与单个计算相关的读取的扭曲边界。要了解如何执行此操作,我建议您查看 this presentation关于基于洗牌的减少。
关于c++ - CUDA:重载共享内存以实现多个数组的缩减方法,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/47896008/
我目前正在学习使用 C 和 OpenMP 进行并行编程。 我想编写简单的代码,其中两个共享值由多个线程递增。 首先我使用了减少指令,它按预期工作。然后我改用 关键 启动关键部分的指令 - 它也有效。
我在用 kubectl scale --replicas=0 -f deployment.yaml 停止我所有正在运行的 pod 。请让我知道是否有更好的方法将所有正在运行的 pod 降到零,保持配置
请考虑我从教程中获得的以下代码和随附的解释性图像。其目的是演示 CUDA 的并行缩减。 #include "cuda_runtime.h" #include "device_launch_parame
我有以下“Frankenstein”和减少代码,部分来自 common CUDA reduction slices ,部分来自 CUDA 示例。 __global__ void reduce
学习openMP // array b #pragma omp parallel for // reduction(&&: b[i])? for (i=2; i<=N; i++
我目前正在使用以下 Reduction 函数通过 CUDA 对数组中的所有元素求和: __global__ void reduceSum(int *input, int *input2, int *i
假设我有两个 numpy 数组,形状为 (d, f) 的 A 和形状为 (d,) 的 I 包含 0..n 中的索引,例如 I = np.array([0, 0, 1, 0, 2, 1]) A = np
我有一个例程,它使用一个循环来计算给定下方粒子表面的粒子的最小高度。此例程尝试随机位置并计算最小高度,然后返回 x, y, z 值,其中 z 是找到的最小高度。 此例程可以与omp parallel
每个视频都有一个有趣的时刻集合,每个时刻代表一个截屏有趣的时间或代表整个标题的时间。请注意,boxarts 和 interestingMoments 数组都位于树中的相同深度。使用 Array.zip
我有一个 ImageIcon,用作打开此 skillsFrame 的按钮。此图像大小为 100x100 像素。正如您在屏幕截图中看到的那样,如果我只是放置图像,它太大了(这是预期的)。 我的问题是关于
我的任务是从 [[a]] 矩阵中获取一列。 一个简单的解决方案是 colFields :: Int -> [[a]] -> [a] colFields n c = map (!! n) c 当减少一级
问题是:如果我输入 hadoop jar MY.jar name_my_class /用户/用户/输入/用户/用户/输出 我需要的所有类都在MY.jar中,为什么我仍然收到错误 java.lang.N
我正在尝试使用 Nodejs、mongoose 和 MongoDB 来进行映射缩减操作。我有一个相当平坦的模式结构,我想获取每个“命名”对象的值/日期对列表。 map 缩减功能显然有问题,但我不知道如
我在 DigitalOcean 中设置了一个 Kubernetes 集群。集群配置为使用 HPA(Horizontal Pod Autoscaler)自动扩展。我想防止终止在过去 1 小时内按比例
IBM Cloudant NoSQL 对每秒的查找、写入、查询有一些限制。 在CloudAnt上我可以编写一个DesignDocument“View”。 当我读取一个 View 时,该读取会对哪里产生
就目前情况而言,这个问题不太适合我们的问答形式。我们希望答案得到事实、引用资料或专业知识的支持,但这个问题可能会引发辩论、争论、民意调查或扩展讨论。如果您觉得这个问题可以改进并可能重新开放,visit
我最近刚刚开始使用 Python 编码,还有很多东西需要学习。我的代码的目标是从单元格中提取字符串,检查其字符长度并用特定缩写替换单词。然后,我将新字符串写入另一个 Excel 工作表中,并在所有数据
我有一个以下形式的 map : Map> START 让 INNER 成为内部映射,即 Map 例如,我想在新 map 中缩小 START map Map END 它们具有相同的键,但具有不同的值。特
给定以下 lambda 表达式,其中 \ 类似于 lambda: (\kf.f(\c.co)km)(\x.dox)(\le.le) 如果我将(\c.co)k转换成ko是不是错了?我这样做了,显然,这是
从 OpenMP 4.0 开始,支持用户定义的缩减。所以我在 C++ 中完全从 here 定义了对 std::vector 的归约。 .它适用于 GNU/5.4.0 和 GNU/6.4.0,但它返回随
我是一名优秀的程序员,十分优秀!