【问题标题】:Use dynamic shared memory allocation for two different vectors对两个不同的向量使用动态共享内存分配
【发布时间】:2013-03-15 14:58:50
【问题描述】:

在内核函数中,我想要两个共享内存向量,长度均为size(实际上是sizeof(float)*size)。

如果需要变量,由于无法直接在内核函数中分配内存,所以我不得不动态分配它,例如:

    myKernel<<<numBlocks, numThreads, 2*sizeof(float)*size>>> (...);  

并且,在内核内部:

extern __shared__ float row[];
extern __shared__ float results[];    

但是,这行不通。

我使用分配的2*size 内存只创建了一个包含所有数据的向量extern __shared__ float rowresults[]。所以row 调用还是一样的,results 调用就像rowresults[size+previousIndex]。这确实有效。

这不是一个大问题,因为无论如何我都得到了预期的结果,但是有没有办法将我动态分配的共享内存分成两个(或更多)不同的变量?只为美丽。

【问题讨论】:

    标签: cuda


    【解决方案1】:

    C Programming guide section on __shared__ 包含从动态分配的共享内存分配多个数组的示例:

    extern __shared__ float array[];
    __device__ void func()      // __device__ or __global__ function
    {
        short* array0 = (short*)array; 
        float* array1 = (float*)&array0[128];
        int*   array2 =   (int*)&array1[64];
    }
    

    由于您只是获取一个指向元素的指针并将其创建为一个新数组,因此我相信您可以调整它以使用动态偏移量而不是示例中的静态偏移量。他们还指出,对齐方式必须相同,这在您的情况下应该不是问题。

    【讨论】:

    • 一种更漂亮的方法,这就是我一直在寻找的。谢谢。
    • @Imortenson 这种方法是否支持内存中元素动态访问对齐的好​​处?如果您已将内存分配为sizeof(float) 或访问将遵循float 对齐方式,那么float4 是否会以32byte 对齐方式访问?谢谢。
    • 在您的示例中,您知道数组的大小;那么为什么不定义一个struct,并将共享内存重新解释为该结构呢?
    • @einpoklum 我复制的示例来自 C 指南,它们是静态偏移量,但提问者事先不知道其数组的大小,因此转换它的用处不大到一个结构。
    • array 应该在func 之外还是在函数内部?
    猜你喜欢
    • 2016-10-21
    • 2012-10-23
    • 1970-01-01
    • 2017-11-11
    • 1970-01-01
    • 2014-06-16
    • 1970-01-01
    • 1970-01-01
    相关资源
    最近更新 更多