常量内存中的推力::device_vector

Posted 2023-02-19

【中文标题】常量内存中的推力::device_vector

【英文标题】：thrust::device_vector in constant memory

【发布时间】：2013-06-08 11:34:57

【问题描述】：

我有一个需要在设备上多次引用的浮点数组，所以我相信存储它的最佳位置是在__常量__内存中（使用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

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

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

【参考方案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);

【讨论】：

感谢您的帮助！向量将是 2 个元素长的幂，可能 >=8096，所以我将放弃使用 __ 常量 __ 内存的想法

如果我更改为全局 device_vector 并引用它，我会在运行时崩溃（嗯，调试运行时）我可以添加全局 device_vector 还是需要在main() 并通过引用传递？

2 的幂或大小不是在这里不使用__constant__ 的原因——正如我所说：你的不是__constant__ 优化的内存访问模式类型。关于你的崩溃：为什么要让它成为全球性的？我看到将其设为全局的问题是您将无法创建在运行时确定的大小的数组，因为构造函数将在 main() 之前调用。跨编译单元构建全局变量的顺序也存在棘手的问题。通常我会在一个函数中创建它并通过引用传递它。

@harrism 请问，你能澄清一下吗？您说：“CUDA 中的常量缓存针对经线中的线程间的统一访问进行了优化。这意味着经线中的所有线程同时访问同一位置。” - 是否意味着如果我使用随机访问__constant__ 内存，那么与cudaMemalloc() 分配的全局内存相比，它不会有任何优势？但是在这种情况下如何加快内存访问，我应该使用LDG load 吗？ on-demand.gputechconf.com/gtc/2013/presentations/…