cuda - 对两个不同的向量使用动态共享内存分配

Question

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

如果需要一个变量，则无法直接在内核函数中分配内存，因此我必须动态分配它，例如:

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

并且，在内核内部:

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

但是，这行不通。

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

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

Answer 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];
}

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