常量内存中的thrust::device_vector

Question

我有一个需要在设备上多次引用的浮点数组，所以我认为存储它的最佳位置是__ constant __ memory (使用this reference)。在初始化时，数组(或向量)需要在运行时写入一次，但由多个不同的函数读取数百万次，因此每次函数调用都不断地复制到内核似乎是一个坏主意。

const int n = 32;
__constant__ float dev_x[n]; //the array in question

struct struct_max : public thrust::unary_function<float,float> {
    float C;
    struct_max(float _C) : C(_C) {}
    __host__ __device__ float operator()(const float& x) const { return fmax(x,C);}
};
void foo(const thrust::host_vector<float> &, const float &);

int main() {
    thrust::host_vector<float> x(n);
    //magic happens populate x
    cudaMemcpyToSymbol(dev_x,x.data(),n*sizeof(float));

    foo(x,0.0);
    return(0);
}

void foo(const thrust::host_vector<float> &input_host_x, const float &x0) {
    thrust::device_vector<float> dev_sol(n);
    thrust::host_vector<float> host_sol(n);

    //this method works fine, but the memory transfer is unacceptable
    thrust::device_vector<float> input_dev_vec(n);
    input_dev_vec = input_host_x; //I want to avoid this
    thrust::transform(input_dev_vec.begin(),input_dev_vec.end(),dev_sol.begin(),struct_max(x0));
    host_sol = dev_sol; //this memory transfer for debugging

    //this method compiles fine, but crashes at runtime
    thrust::device_ptr<float> dev_ptr = thrust::device_pointer_cast(dev_x);
    thrust::transform(dev_ptr,dev_ptr+n,dev_sol.begin(),struct_max(x0));
    host_sol = dev_sol; //this line crashes
}

我尝试添加一个全局thrust::device_vector dev_x(n)，但它在运行时也崩溃了，并且可能在__全局__内存而不是__ constant__内存中

如果我只需要丢弃推力库，这一切都可以工作，但是有没有办法使用带有全局变量和设备常量内存的推力库？

Answer 1

问得好！您不能将__constant__数组转换为常规设备指针。

我将回答您的问题(在下面一行之后)，但首先:这是对__constant__的错误使用，它并不是您真正想要的。CUDA中的常量缓存针对warp中跨线程的统一访问进行了优化。这意味着扭曲中的所有线程同时访问相同的位置。如果warp的每个线程访问不同的常量内存位置，那么这些访问将被序列化。因此，您的访问模式，即连续线程访问连续内存位置，将比统一访问慢32倍。你真的应该只使用设备内存。如果您需要一次写入数据，但需要多次读取，那么只需使用device_vector:初始化它一次，然后读取它多次。

要完成所要求的操作，您可以使用thrust::counting_iterator作为thrust::transform的输入，以便在__constant__数组中生成一系列索引。然后您的函数器的operator()接受int索引操作数而不是float值操作数，并在常量内存中进行查找。

(请注意，这意味着您的函数器现在只是__device__代码。如果您需要可移植性，您可以很容易地重载操作符以获取一个浮点数，并在主机数据上以不同的方式调用它。)

我修改了您的示例，以初始化数据并打印结果以验证其正确性。

#include <stdio.h>
#include <stdlib.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/iterator/counting_iterator.h>

const int n = 32;
__constant__ float dev_x[n]; //the array in question

struct struct_max : public thrust::unary_function<float,float> {
    float C;
    struct_max(float _C) : C(_C) {}

    // only works as a device function
    __device__ float operator()(const int& i) const { 
        // use index into constant array
        return fmax(dev_x[i],C); 
    }
};

void foo(const thrust::host_vector<float> &input_host_x, const float &x0) {
    thrust::device_vector<float> dev_sol(n);
    thrust::host_vector<float> host_sol(n);

    thrust::device_ptr<float> dev_ptr = thrust::device_pointer_cast(dev_x);
    thrust::transform(thrust::make_counting_iterator(0),
                      thrust::make_counting_iterator(n),
                      dev_sol.begin(),
                      struct_max(x0));
    host_sol = dev_sol; //this line crashes

    for (int i = 0; i < n; i++)
        printf("%f\n", host_sol[i]);
}

int main() {
    thrust::host_vector<float> x(n);

    //magic happens populate x
    for (int i = 0; i < n; i++) x[i] = rand() / (float)RAND_MAX;

    cudaMemcpyToSymbol(dev_x,x.data(),n*sizeof(float));

    foo(x, 0.5);
    return(0);
}