logo标签列表

CUDA Thrust使用double2数组进行约简

cudagpgputhrust
4

我有下面这段（可编译且可执行）使用CUDA Thrust来对float2数组进行归约的代码。它能正常工作。

using namespace std;

// includes, system 
#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>
#include <conio.h>

#include <typeinfo>  
#include <iostream>

// includes CUDA
#include <cuda.h>
#include <cuda_runtime.h>

// includes Thrust
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/reduce.h>

// float2 + struct
struct add_float2 {
    __device__ float2 operator()(const float2& a, const float2& b) const {
        float2 r;
        r.x = a.x + b.x;
        r.y = a.y + b.y;
        return r;
    }
 };

// double2 + struct
struct add_double2 {
    __device__ double2 operator()(const double2& a, const double2& b) const {
        double2 r;
        r.x = a.x + b.x;
        r.y = a.y + b.y;
        return r;
    }
 };

void main( int argc, char** argv) 
{
    int N = 20;

    // --- Host
    float2* ha; ha = (float2*) malloc(N*sizeof(float2));
    for (unsigned i=0; i<N; ++i) {
        ha[i].x = 1;
        ha[i].y = 2;
    }

    // --- Device
    float2* da; cudaMalloc((void**)&da,N*sizeof(float2));
    cudaMemcpy(da,ha,N*sizeof(float2),cudaMemcpyHostToDevice);

    thrust::device_ptr<float2> dev_ptr_1(da);
    thrust::device_ptr<float2> dev_ptr_2(da+N);

    float2 init; init.x = init.y = 0.0f;

    float2 sum = thrust::reduce(dev_ptr_1,dev_ptr_2,init,add_float2());

    cout << " Real part = " << sum.x << "; Imaginary part = " << sum.y << endl;

    getch();

 }

然而，当我在main程序中将float2更改为double2时，即
void main( int argc, char** argv) 
{
    int N = 20;

    // --- Host
    double2* ha; ha = (double2*) malloc(N*sizeof(double2));
    for (unsigned i=0; i<N; ++i) {
        ha[i].x = 1;
        ha[i].y = 2;
    }

    // --- Device
    double2* da; cudaMalloc((void**)&da,N*sizeof(double2));
    cudaMemcpy(da,ha,N*sizeof(double2),cudaMemcpyHostToDevice);

    thrust::device_ptr<double2> dev_ptr_1(da);
    thrust::device_ptr<double2> dev_ptr_2(da+N);

    double2 init; init.x = init.y = 0.0;

    double2 sum = thrust::reduce(dev_ptr_1,dev_ptr_2,init,add_double2());

    cout << " Real part = " << sum.x << "; Imaginary part = " << sum.y << endl;

    getch();

}

我在reduce行收到一个exception。我如何使用CUDA Thrust对double2数组进行缩减？我有做错什么吗？提前感谢。

根据TALONMIES的答案，以下是可行的解决方案：

using namespace std;

// includes, system
#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>
#include <conio.h>

#include <typeinfo>
#include <iostream>

// includes CUDA
#include <cuda.h>
#include <cuda_runtime.h>

// includes Thrust
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/reduce.h>

struct my_double2 {
    double x, y;
};

// double2 + struct
struct add_my_double2 {
    __device__ my_double2 operator()(const my_double2& a, const my_double2& b) const {
        my_double2 r;
        r.x = a.x + b.x;
        r.y = a.y + b.y;
        return r;
    }
};

void main( int argc, char** argv) 
{
    int N = 20;

    // --- Host
    my_double2* ha; ha = (my_double2*) malloc(N*sizeof(my_double2));
    for (unsigned i=0; i<N; ++i) {
        ha[i].x = 1;
        ha[i].y = 2;
    }

    // --- Device
    my_double2* da; cudaMalloc((void**)&da,N*sizeof(my_double2));
    cudaMemcpy(da,ha,N*sizeof(my_double2),cudaMemcpyHostToDevice);

    thrust::device_ptr<my_double2> dev_ptr_1(da);
    thrust::device_ptr<my_double2> dev_ptr_2(da+N);

    my_double2 init; init.x = init.y = 0.0;

    cout << "here3\n";
    my_double2 sum = thrust::reduce(dev_ptr_1,dev_ptr_2,init,add_my_double2());

    cout << " Real part = " << sum.x << "; Imaginary part = " << sum.y << endl;

    getch();

}
1个回答
4
这是MSVC和nvcc之间已知的不兼容性问题。例如，请参见这里。解决方案是定义您自己版本的double2并使用它代替原来的版本。
仅供参考，我可以在安装了CUDA 5.5的Linux 64位计算机上正确编译和运行您的代码。
非常感谢您的回答。这种不兼容的原因是什么？它是否也涉及到最新版本的MSVC？在引用的网页中，建议将自定义版本的double2定义为struct。性能如何？过去我曾尝试过使用我的double2版本作为一对double数字，但我意识到内部使用double2的包装器类的解决方案比一对double更快。 - Vitality
1
我对微软的编译器和库不是很熟悉，所以我不能告诉你更多关于我在答案中提供的内容。至于double2与结构体之间的区别 - 我想你可能不知道double2实际上就是一个简单的结构体（请查看vector_types.h）。使用CUDA向量类型和使用自己的功能等效结构体在生成的代码或性能上应该没有任何区别（但根据我的经验并非如此）。 - talonmies
谢谢，我接受了你的答案。是的，我知道double2是一个struct，但在定义前面有一些特殊的关键字：__device_builtin__ __builtin_align__(16)。据我所知，它们应该允许协同内存访问。相比不使用对齐的自定义定义，这样做会带来一些改进吗？ - Vitality
2
__device__builtin__只是一个标签，它没有任何作用，而__builtin__align__解析为基本的__align__宏。CUDA“builin”向量类型只是简单的结构体，它们依赖于编译器分析来实现合并，就像任何用户定义的类型一样。它们真正有意义的地方只在于与纹理硬件的关系，即使如此，也仅仅是为了让类型大小和对齐方式与纹理API所期望的相匹配。 - talonmies
网页内容由stack overflow 提供, 点击上面的
可以查看英文原文，
原文链接