Переменная массива Cuda Shared Memory

Question

Я пытаюсь объявить переменную для умножения матриц следующим образом:

__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];

Я пытаюсь сделать так, чтобы пользователь мог ввести размер матрицы для вычисления, однако это означало бы изменениеРАЗМЕР БЛОКА.Я изменил его, но получаю ошибку компилятора: «ошибка: постоянное значение не известно».Я смотрел на это, и это похоже на эту тему .Итак, я попытался:

__shared__ int buf [];

Но затем я получаю: «ошибка: недопустимый тип не разрешен»

Спасибо, Дэн Обновление с кодом (в значительной степени следовало это руководство и просмотр с руководством гида): Размер блока передается, спрашивая пользователя о размере матрицы.Они вводят х и у.Размер блока равен только 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;


}

Answer 1

extern __shared__ int buf[];

когда вы запускаете ядро, вы должны запустить его таким образом;

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

Если у вас есть несколько внешних объявлений общего доступа:

extern __shared__ float As[];

extern __shared__ float Bs[];

это приведет к As, указывающему на тот же адрес, что и Bs.

Вам необходимо сохранитьКак и B внутри 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]

Answer 2

У вас есть два варианта объявления общей памяти внутри ядра - статическая или динамическая. Я предполагаю, что то, что вы делаете в данный момент, выглядит примерно так:

#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 >>>(....);

Все они одинаково действительны, хотя я лично предпочитаю версию шаблона, потому что она может позволить другую оптимизацию компилятора, такую ​​как автоматическое развертывание цикла, что динамическая версия не может без дополнительной работы.

Answer 3

Звучит правильно.

Обычно в этом случае вам нужно что-то неправильно распределить.

Здесь есть две вещи, одна Си не знает о 2D-массивах (это просто массив массивов), и размеры массивов должны компилировать постоянные времени (или что-то, что компилятор может вычислить во время компиляции).

Если вы используете C99, вы можете объявить размер массива с помощью параметра функции, но поддержка C99 ... в лучшем случае пятнистая.