标签: reduction

我需要执行并行缩减以在CUDA设备上找到数组的最小值或最大值.我找到了一个很好的图书馆,称为Thrust.您似乎只能在主机内存中对阵列执行并行缩减.我的数据在设备内存中.是否可以减少设备内存中的数据？我无法想象如何做到这一点.以下是Thrust的文档:http://code.google.com/p/thrust/wiki/QuickStartGuide#Reductions.谢谢大家.

如何在Delphi中将颜色减少到指定的数字(<= 256)？我不想只使用:

 Bmp.PixelFormat := pf8bit;

因为那样我无法控制颜色数量.我不想要抖动,因为我已经知道如何用256或更少颜色抖动图像.

我发现了这个Median Cut实现,但它从1990年开始是纯粹的Pascal,并且:

1. 不在Delphi中编译
2. 说它是共享软件并且花费25 Deutche Marks
3. 看起来(不知何故)不必要的复杂

我想TBitmap32只减少(Graphics32位图类,仅支持32位颜色)到<= 8bit bmp.我不需要减少到15/16bit,我不需要从24或15/16bit图像减少.只需32位=> 8位 -

Delphi我使用:7,2005,XE3.

OpenMP标准指定了缩减变量的初始值.所以我必须初始化变量,在下列情况下我该怎么做:

int sum;
//...
for(int it=0;i<maxIt;i++){
#pragma omp parallel
{
  #pragma omp for nowait
  for(int i=0;i<ct;i++)
    arrayX[i]=arrayY[i];

  sum = 0;
  #pragma omp for reduction(+:sum)
  for(int i=0;i<ct;i++)
    sum+=arrayZ[i];
}
//Use sum
}

请注意,我只使用1个并行区域来最小化开销并允许第一个循环中的nowait.使用这个原样将导致数据竞争(IMO),因为在其他线程启动第二个循环之后来自第一个循环的线程将重置总和.
当然,我可以在外部循环的顶部执行此操作,但在一般情况下,对于大型代码库,您可能会忘记您需要或已将其设置在那里会产生意外结果.
"omp single"有帮助吗？我怀疑当线程A执行单个时,另一个线程可能已进入减少循环."omp障碍"是可能的,但我想避免它,因为它击败了"nowait".

又一个例子:

#pragma omp parallel
{
  sum = 0;
  #pragma omp for reduction(+:sum)
  for(int i=0;i<ct;i++)
    sum+=arrayZ[i];
  //Use sum
  sum = 0;
  #pragma omp for reduction(+:sum)
  for(int i=0;i<ct;i++)
    sum+=arrayZ[i];
  //Use sum
}

我将如何(重新)初始化？

对于大小为60000行,10列的2D数组,我有如下阵列

 [0 0 0 0 0 1 0 0 0 0]
 [1 0 0 0 0 0 0 0 0 0]
 [0 0 0 0 0 0 0 0 0 1]
.......

任何行只包含一个'1'

我必须将它减少到一个行或列向量,它显示我们有一个1的索引.例如,对于上面显示的行,我们必须最终得到

[6,1,10...] 超过第60,000个值.

如何在没有循环的Matlab中做到这一点？

前一个问题询问如何有效地找到CUDA中数组的最大值:在CUDA中查找最大值,最高响应提供了一个关于优化还原内核的NVIDIA演示文稿的链接.

如果您使用的是Visual Studio,只需删除标题引用,以及CPU EXECUTION之间的所有内容.

我设置了一个找到max的变体,但它与CPU发现的不匹配:

// Returns the maximum value of
// an array of size n
float GetMax(float *maxes, int n)
{
    int i = 0;
    float max = -100000;
    for(i = 0; i < n; i++)
    {
        if(maxes[i] > max)
            max = maxes[i];
    }

    return max;
}

// Too obvious...
__device__ float MaxOf2(float a, float b)
{
    if(a > b)   return a;
    else            return b;
}


__global__ void MaxReduction(int n, float *g_idata, float *g_odata) …

If I use a barrier (no matter if CLK_LOCAL_MEM_FENCE or CLK_GLOBAL_MEM_FENCE) in my kernel, it causes a CL_INVALID_WORK_GROUP_SIZE error. The global work size is 512, the local work size is 128, 65536 items have to be computed, the max work group size of my device is 1024, I am using only one dimension. For Java bindings I use JOCL. The kernel is very simple:

kernel void sum(global float *input, global float *output, const int numElements, local float *localCopy
{ …

我正在尝试沿着2D矩阵的行方向实现缩减.我从stackoverflow上找到的代码开始(非常感谢Robert!)

thrust :: max_element比较cublasIsamax慢 - 更有效的实现？

上面的链接显示了一个在单行上执行缩减的自定义内核.它将输入行分为多行,每行有1024个线程.效果很好.

对于2D情况,一切都是相同的,除了现在有一个网格尺寸.所以每个块的y维度仍然是1.问题是当我尝试将数据写入每个块内的共享内存(在代码中的"max_idx_kernel_reduction_within_block"内核中)时,需要很长时间(超过(行数)*(在1行上执行减少所需的时间.我宁愿运行for循环).我知道我有很多元素,但我期待比这更快的东西.

我不认为内存访问模式是一个问题,但我听说TOTAL共享内存量可能是限制？:CUDA:合并全局内存访问速度比共享内存快吗？另外,分配大型共享内存阵列会减慢程序的速度吗？

任何使我的代码更快的建议(第一个内核是瓶颈)？非常感谢,非常感谢!!

#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <iostream>
#include <cuda_runtime.h>

#define NCOLS 163317 // number of columns
#define NROWS 8 // number of rows
#define nTPB 1024  // Threads per Block. nTPB should be a power-of-2
#define MAX_BLOCKS_X ((NCOLS/nTPB)+1) // # of blocks I will launch

#define MIN(a,b) ((a>b)?b:a)
#define FLOAT_MIN -1.0f // lowest anticipated number of the data. Values in array will be compared with this and updated …

我正在 CUDA 中实现并行缩减。

内核等待__syncthreads所有线程完成对共享内存的 2 次读取，然后将总和写回共享内存。

我应该使用 a__threadfence_block来确保下一次迭代的所有线程都可以看到对共享内存的写入，还是__syncthreads按照NVIDIA 示例中给出的方式使用？

我对 CUDA 编程相当陌生，我正在尝试编写一个 CUDA 内核，用于仅对 3 维张量的 1 维进行并行缩减，该张量是float馈入内核的行主展平数组。

换句话说，我试图用,和numpy.sum的有限轴组合重写。axis=0axis=1axis=2

我已经成功实现了“reduce over axis 0”和“reduce over axis 1”，但是“reduce over axis2”的性能问题让我在这里发布了一个问题来寻求建议。

内核以一维网格和一维块配置启动，并将每个线程映射到缩减输出张量的每个元素。所以，它应该是这样的：

这是我的内核：

__global__ void kernel_reduce_sum_3d_try02(
        float* g_idata,
        float* g_odata,
        int dim0,
        int dim1,
        int dim2,
        int overaxis0,
        int overaxis1,
        int overaxis2)
{

    if (overaxis0 == 0 && overaxis1 == 0 && overaxis2 == 1) { 
        // static shared memory
        __shared__ float smem_store[BLOCK_SIZE];

        // set thread ID
        //unsigned int tid = threadIdx.x;
        unsigned int tid …

我正在尝试使用 FFT 与使用加窗方法比较互相关。

我的 Matlab 代码是：

isize = 20;
n = 7;
for i = 1:n %%7x7 xcorr
  for j = 1:n
    xcout(i,j) = sum(sum(ffcorr1 .* ref(i:i+isize-1,j:j+isize-1))); %%ref is 676 element array and ffcorr1 is a 400 element array
  end
end

类似的CUDA内核：

__global__ void xc_corr(double* in_im, double* ref_im, int pix3, int isize, int n, double* out1, double* temp1, double* sum_temp1)
{

    int p = blockIdx.x * blockDim.x + threadIdx.x;
    int q = 0;
    int i = 0;
    int j = …

我想计算Cuda中整个图像的平均值.为了测试2D数组的减少效果,我在下面编写了这个内核.最终输出o​​应该是所有图像值的总和.输入g是2D阵列,每个像素的值为1.但是这个程序的结果是总和为0.对我来说有点奇怪.

我在本教程中模仿1D阵列的减少http://developer.download.nvidia.com/compute/cuda/1.1-Beta/x86_website/projects/reduction/doc/reduction.pdf我写这个2D表格.我是Cuda的新手.欢迎提出有关潜在错误和改进的建议!

只需添加一条评论.我知道计算一维数组的平均值是有意义的.但我想利用更多并测试更复杂的还原行为.这可能不对.但只是一个考验.希望任何人都能给我更多关于减少常见做法的建议.

#include <iostream>
#include <stdio.h>
#include <stdlib.h>
#include <time.h>

cudaEvent_t start, stop;
float elapsedTime;

__global__ void 
reduce(float *g, float *o, const int dimx, const int dimy)
{
extern __shared__ float sdata[];

unsigned int tid_x = threadIdx.x;
unsigned int tid_y = threadIdx.y;

unsigned int i = blockDim.x * blockIdx.x + threadIdx.x;
unsigned int j = blockDim.y * blockIdx.y + threadIdx.y; 

if (i >= dimx || j >= dimy)
    return;

sdata[tid_x*blockDim.y + tid_y] = g[i*dimy + j];

__syncthreads(); …