我尝试为矩阵乘法声明一个变量,如下所示:
__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
我正在尝试使它,以便用户可以输入矩阵的大小来计算,但这将意味着更改BLOCK_SIZE。我更改了它,但我得到了一个编译器错误:“错误:常量值未知”。我已经研究过了,它和这个thread很相似。所以我试着:
__shared__ int buf [];
但是我得到的是:“错误:不允许不完整的类型”
谢谢,Dan Update with code(基本上遵循了this guide和staring with cuda指南):块大小是通过询问用户矩阵的大小传入的。他们输入x和y。块大小只有x,现在它必须接受与x和y相同的大小。
__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;
}
发布于 2012-02-08 17:50:29
extern __shared__ int buf[];
当你启动内核时,你应该这样启动它;
kernel<<<blocks,threads,numbytes_for_shared>>>(...);
如果您有多个shared外部声明:
extern __shared__ float As[];
extern __shared__ float Bs[];
这将导致As
指向与Bs
相同的地址。
您将需要将As和Bs保持在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]
发布于 2012-02-08 17:42:27
在内核中声明共享内存有两种选择--静态或动态。我假设你现在正在做的事情是这样的:
#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 >>>(....);
所有这些都是同样有效的,尽管我个人更喜欢模板版本,因为它可以支持其他编译器优化,如自动循环展开,而动态版本在没有额外工作的情况下无法实现。
发布于 2012-02-08 16:26:54
听起来是对的。
通常情况下,在这种情况下,您需要malloc一些东西。
这里有两件事,一是C不知道二维数组(它只是一个数组的数组)和数组大小需要编译时间常量(或者编译器可以在编译时计算的东西)。
如果您使用的是C99,则可以使用函数的参数声明数组大小,但C99支持...充其量也就是斑点。
https://stackoverflow.com/questions/9187899
复制相似问题