Cuda 共享内存数组变量

18

我正在尝试声明一个矩阵乘法的变量,如下所示:

__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];

我正在尝试让用户输入矩阵的大小以进行计算,但这意味着需要更改BLOCK_SIZE。我已经更改了它,但是我遇到了编译器错误:

error: constant value is not known

我已经研究过了,它类似于这个thread。所以我尝试了:

__shared__ int buf [];

但是我得到:

error: incomplete type is not allowed

谢谢,丹

更新代码(基本上遵循this guide和CUDA指南的开始):
块大小是通过询问矩阵大小的用户传递的。他们输入xy。块大小仅为x,现在必须接受与xy相同的大小。

__global__ void matrixMul( float* C, float* A, float* B, int wA, int wB,size_t block_size)
{
    // Block index
    int bx = blockIdx.x;
    int by = blockIdx.y;
    
    // Thread index
    int tx = threadIdx.x;
    int ty = threadIdx.y;
    
    // Index of the first sub-matrix of A processed 
    // by the block
    int aBegin = wA * block_size * by;
 
    // Index of the last sub-matrix of A processed 
    // by the block
    int aEnd   = aBegin + wA - 1;
 
    // Step size used to iterate through the 
    // sub-matrices of A
    int aStep  = block_size;
 
    // Index of the first sub-matrix of B processed 
    // by the block
    int bBegin = block_size * bx;
 
    // Step size used to iterate through the 
    // sub-matrices of B
    int bStep  = block_size * wB;
    float Csub=0;
    // Loop over all the sub-matrices of A and B
    // required to compute the block sub-matrix
    for (int a = aBegin, b = bBegin; a <= aEnd; a += aStep, b += bStep) 
    {
        // Declaration of the shared memory array As 
        // used to store the sub-matrix of A
        
        extern __shared__ float As[];

        // Declaration of the shared memory array Bs 
        // used to store the sub-matrix of B
        extern __shared__ float Bs[];
        extern __shared__ float smem[];

        // Load the matrices from global memory
        // to shared memory; each thread loads
        // one element of each matrix
        smem[ty*block_size+tx] = A[a + wA * ty + tx];
        //cuPrintf("\n\nWhat are the memory locations?\n");
        //cuPrintf("The shared memory(A) is: %.2f\n",smem[ty*block_size+tx]);
        smem[block_size*block_size+ty*block_size+tx]  = B[b + wB * ty + tx];
        //cuPrintf("The shared memory(B) is: %.2f\n",smem[block_size*block_size+ty*block_size+tx]);
        // Synchronize to make sure the matrices 
        // are loaded
        __syncthreads();
 
        // Multiply the two matrices together;
        // each thread computes one element
        // of the block sub-matrix
        for (int k = 0; k < block_size; ++k)
        {
        
            Csub += smem[ty*block_size+k] * smem[block_size*block_size+k*block_size+tx] ;
            //cuPrintf("Csub is currently: %.2f\n",Csub);
        }
        //cuPrintf("\n\n\n");
        // Synchronize to make sure that the preceding
        // computation is done before loading two new
        // sub-matrices of A and B in the next iteration
        //cuPrintf("the results are csub: %.2f\n",Csub);
        __syncthreads();
    }
    // Write the block sub-matrix to device memory;
    // each thread writes one element
    int c = wB * block_size * by + block_size * bx;
    C[c + wB * ty + tx] = Csub;
    
    
}
3个回答

34
extern __shared__ int buf[];

当你启动内核时,应该以这种方式启动它;

kernel<<<blocks,threads,numbytes_for_shared>>>(...);

如果您有多个共享变量的外部声明:

extern __shared__ float As[];
// ...
extern __shared__ float Bs[];

这将导致As指向与Bs相同的地址。

您需要将AsBs保留在1D数组中。

extern __shared__ float smem[];

在调用内核时,您应该使用2*BLOCK_SIZE*BLOCK_SIZE*sizeof(float)来启动它。

当索引到As时,请使用smem[y*BLOCK_SIZE+x],当索引到Bs时,请使用smem[BLOCK_SIZE*BLOCK_SIZE+y*BLOCK_SIZE+x]


谢谢提供的信息。我现在只是在矩阵乘法方面遇到了一点问题,但会根据这里提供的信息尝试解决。 - Dan
@brano 这个方法适用于4x4矩阵,但是一旦超过这个大小,矩阵C的结果就会出现问题。数值要么全是0,要么混合着0。 - Dan
@brano,我已经更新了我的函数代码以反映我所做的更改。 - Dan
@Dan:内核没有问题。我怀疑你以错误的方式启动了内核。矩阵维度需要是块大小的倍数。此外,我在共享内存的大小上犯了一个错误。它应该是字节数(2 * block_size * block_size * sizeof(float))。 - brano
@brano 谢谢,问题已经解决了。但是它又产生了另一个问题,我正在研究中。也许我会就那个问题发布另一个问题。感谢你和talonmies的帮助。 - Dan
显示剩余3条评论

32

在内核中声明共享内存有两种选择 - 静态或动态。我假设你目前正在做的事情类似于这样:

#define BLOCK_SIZE (16)

__global__ void sgemm0(const float *A, const float *B, float *C)
{
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];

}

你希望能够轻松地更改BLOCK_SIZE。

一种可能的方法是继续使用静态共享内存分配,但将分配大小作为模板参数,像这样:

template<int blocksize=16>
__global__ void sgemm1(const float *A, const float *B, float *C)
{
    __shared__ float As[blocksize][blocksize];

}
template void sgemm1<16>(const float *, const float *, float *C);

那么您可以在编译时实例化所需的许多不同块大小变体。

如果您想要动态分配内存,可以这样定义:

__global__ void sgemm2(const float *A, const float *B, float *C)
{
    extern __shared__ float As[];

} 

然后将分配的大小作为参数添加到内核调用中:

size_t blocksize = BLOCK_SIZE * BLOCK_SIZE;
sgemm2<<< gridDim, blockDim, sizeof(float)*blocksize >>>(....);

如果你有多个静态声明的数组,希望用动态分配的共享内存替换它们,请注意每个核函数只有一个动态共享内存分配,因此多个元素存在于该内存段中。所以如果你有类似以下的内容:

#define BLOCK_SIZE (16)

__global__ void sgemm0(const float *A, const float *B, float *C)
{
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];

}

您可以将其替换为:

#define BLOCK_SIZE (16)

__global__ void sgemm3(const float *A, const float *B, float *C)
{
    extern __shared__ float buffer[];

    float *As = &buffer[0];
    float *Bs = &buffer[BLOCK_SIZE*BLOCK_SIZE];

}

然后像这样启动内核:

size_t blocksize = 2 * BLOCK_SIZE * BLOCK_SIZE;
sgemm3<<< gridDim, blockDim, sizeof(float)*blocksize >>>(....);

所有方法都是有效的,尽管我个人更喜欢使用模板版本,因为它可以允许其他编译器优化,如自动循环展开,而动态版本没有这个功能需要额外的工作。


1
@talonmies,那个__shared__ float *As; 应该改为extern shared float As[];,就像brano的回答一样。你们两个都会得到点赞。 :) - harrism
啊,错过了那个。谢谢马克。 - talonmies
好的,谢谢你的回答。我已经使用了extern shared。但是它将其转换为一维数组,程序最初使用的是二维数组。我知道二维数组是数组的数组,所以一维数组应该可以工作。例如:As[ty][tx] = A[a + wA * ty + tx]; 我将其转换为 -> As[ty*MAX_THREADS+tx] = A[a + wA * ty + tx]; 其中maxthreads为1023,因为我的最大线程数为1024。但是对于一个2x2矩阵,我得到了-0,0,-0,0。 - Dan
我也尝试了那个。实际上,我得到了结果,但它们是不正确的。我正在更新我的原始帖子以反映我现在使用的代码。 - Dan
默认模板参数方法从未出现在我的脑海中。我想说“太棒了!”和“谢谢!”但没有声音。 - 3Dave
显示剩余3条评论

0

听起来没问题。

通常在这种情况下,您需要分配一些内存。

这里有两件事情,一是C不知道2D数组(它只是一个数组的数组),二是数组大小需要编译时常量(或者编译器可以在编译时计算的东西)。

如果您正在使用C99,则可以使用函数的参数声明数组大小,但是C99支持最好是...不稳定的。


我尝试过使用malloc,但我不认为在设备代码上允许这样做。 - Dan
无法在设备代码上调用 malloc ... 所有动态内存都必须在进入内核之前分配,动态缓冲区需要使用CUDA特定版本的 mallocmemcpy 在设备上分配和复制。 - Jason
@Jason:实际上,在Fermi GPU上,malloc和C++的new运算符都被支持。但仅适用于将驻留在全局内存中的分配。您正确地断言动态分配的共享内存必须由调用主机代码(在这种情况下作为内核启动语法的一部分或通过单独的API调用)进行分配。 - talonmies
@talonmies:那么在设备和主机之间分配和释放内存时,您不再需要使用“cudaMalloc()”、“cudaMemcpy()”和“cudaFree()”的序列了吗?如果是这样,那是在哪个版本的Cuda中更改的? - Jason
@Jason,通常这仍然是最好的工作方式,但线程可以根据需要从运行时堆中分配自己的全局内存。Cuda 3.1引入了内核malloc支持,CUDA 4.0添加了new运算符。它仅在计算能力为2.0和2.1的设备上受支持,目前性能并不特别好,但是它是被支持的。 - talonmies

网页内容由stack overflow 提供, 点击上面的
可以查看英文原文,
原文链接