2

我有两个向量（oldvector和newvector）。我需要计算由以下伪代码定义的残差值：

residual = 0;
forall i : residual += (oldvector[i] - newvector[i])^2

目前，我正在使用两个 CUDA Thrust 操作来计算这个，这两个操作本质上是在做的：

forall i : oldvector[i] = oldvector[i] - newvector[i]

后跟thrust::transform_reduce一个正方形作为一元运算符，它正在做：

residual = 0;
forall i : residual += oldvector[i]^2;

这个问题显然是之前全局内存的中间存储transform_reduce。有没有更有效的方法来解决这个问题，可以融合这两个步骤？除了编写自己的 CUDA 内核，还有其他选择吗？

我想到的一种方法是thrust::reduce用 zip 迭代器编写一个。这样做的问题是运算符的返回类型必须与其输入的类型相同。在我看来，这意味着归约运算符将返回一个元组，这意味着额外的加法。

如果我确实编写了缩减 CUDA 内核，那么缩减内核的 CUDA 1.1 示例是否有任何改进？

4

2 回答 2

3

推力::inner_product将在单个函数调用中完成。也可以使您的原始想法起作用（将两个向量压缩在一起并使用thrust::transform_reduce）此代码演示了这两种方法：

#include <iostream>

#include <thrust/tuple.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/transform.h>
#include <thrust/device_vector.h>
#include <thrust/inner_product.h>
#include <thrust/functional.h>

#define N 2

struct zdiffsq{
template <typename Tuple>
  __host__ __device__ float operator()(Tuple a)
  {
    float result = thrust::get<0>(a) - thrust::get<1>(a);
    return result*result;
  }
};

struct diffsq{
  __host__ __device__ float operator()(float a, float b)
  {
    return (b-a)*(b-a);
  }
};

int main(){

  thrust::device_vector<float> oldvector(N);
  thrust::device_vector<float> newvector(N);
  oldvector[0] = 1.0f;  oldvector[1] = 2.0f;
  newvector[0] = 2.0f;  newvector[1] = 5.0f;

  float result = thrust::inner_product(oldvector.begin(), oldvector.end(), newvector.begin(), 0.0f, thrust::plus<float>(), diffsq());
  std::cout << "Result: " << result << std::endl;

  float result2 = thrust::transform_reduce(thrust::make_zip_iterator(thrust::make_tuple(oldvector.begin(), newvector.begin())), thrust::make_zip_iterator(thrust::make_tuple(oldvector.end(), newvector.end())), zdiffsq(), 0.0f, thrust::plus<float>());
  std::cout << "Result2: " << result2 << std::endl;
}

您还可以通过使用推力占位符来研究消除与内积示例一起使用的函子定义。

即使您想编写自己的 CUDA 代码，现在对于常用算法（如并行归约和排序）的标准推荐是使用cub。

是的，CUDA 并行缩减示例和随附的演示文稿仍然是快速并行缩减的一个很好的基本介绍。

于 2014-05-11T23:17:48.043 回答
1

Robert Crovella 已经回答了这个问题，并且还建议使用CUB。

与 Thrust 不同的是，CUB 保留了性能关键参数，例如要使用的特定缩减算法的选择和未绑定的并发程度，由用户选择。可以调整这些参数以最大限度地提高特定架构和应用程序的性能。可以在编译时指定参数，从而避免运行时性能损失。

下面，有一个关于如何使用 CUB 进行残差计算的完整示例。

#include <cub/cub.cuh>
#include <cuda.h>

#include "Utilities.cuh"

#include <iostream>

#define BLOCKSIZE           256
#define ITEMS_PER_THREAD    8

const int N = 4096;

/******************************/
/* TRANSFORM REDUCTION KERNEL */
/******************************/
__global__ void TransformSumKernel(const float * __restrict__ indata1, const float * __restrict__ indata2, float * __restrict__ outdata) {

    unsigned int tid = threadIdx.x + blockIdx.x * gridDim.x;

    // --- Specialize BlockReduce for type float. 
    typedef cub::BlockReduce<float, BLOCKSIZE * ITEMS_PER_THREAD> BlockReduceT;

    __shared__ typename BlockReduceT::TempStorage  temp_storage;

    float result;
    if(tid < N) result = BlockReduceT(temp_storage).Sum((indata1[tid] - indata2[tid]) * (indata1[tid] - indata2[tid]));

    if(threadIdx.x == 0) atomicAdd(outdata, result);

    return;
}

/********/
/* MAIN */
/********/
int main() {

    // --- Allocate host side space for 
    float *h_data1      = (float *)malloc(N * sizeof(float));
    float *h_data2      = (float *)malloc(N * sizeof(float));
    float *h_result     = (float *)malloc(sizeof(float));

    float *d_data1;     gpuErrchk(cudaMalloc(&d_data1, N * sizeof(float)));
    float *d_data2;     gpuErrchk(cudaMalloc(&d_data2, N * sizeof(float)));
    float *d_result;    gpuErrchk(cudaMalloc(&d_result, sizeof(float)));

    for (int i = 0; i < N; i++) {
        h_data1[i] = 1.f;
        h_data2[i] = 3.f;
    }

    gpuErrchk(cudaMemcpy(d_data1, h_data1, N * sizeof(float), cudaMemcpyHostToDevice));
    gpuErrchk(cudaMemcpy(d_data2, h_data2, N * sizeof(float), cudaMemcpyHostToDevice));

    TransformSumKernel<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(d_data1, d_data2, d_result);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());

    gpuErrchk(cudaMemcpy(h_result, d_result, sizeof(float), cudaMemcpyDeviceToHost));

    std::cout << "output: ";
    std::cout << h_result[0];
    std::cout << std::endl;

    gpuErrchk(cudaFree(d_data1));
    gpuErrchk(cudaFree(d_data2));
    gpuErrchk(cudaFree(d_result));

    return 0;
}
于 2015-08-04T07:51:12.497 回答