cuda - 動的並列処理 CUDA での共有メモリの使用

Question

質問 1: 共有メモリが子カーネルによってのみ使用される場合、親カーネルの起動時に割り当てられる動的共有メモリの量を指定する必要がありますか?

質問 2: 以下は私の子カーネルと親カーネルです

親カーネル

__global__ void ColumnFractionalShift(DataIn DataInput,float* __restrict__ DeviceInput, float ShiftAmount, float* __restrict__ LightFieldDevice)
{
    cudaError_t status;
    float ImageShift = threadIdx.x*ShiftAmount;
    float ImageIntegerShift = nearbyintf(ImageShift);
    float Delay = ImageShift - ImageIntegerShift;
    int InputImageOffset = +DataInput.X*DataInput.Y*DataInput.U*(threadIdx.y) + DataInput.X*DataInput.Y*(threadIdx.x);
    dim3 dimBlock(32, 24);
    dim3 dimGrid(16, 14);
    //if (threadIdx.x > 5)
    {
        ConvolutionColumn << <dimGrid, dimBlock, ((sizeof(float)* 24 * 32 * 3)) >> >(DataInput, DeviceInput + InputImageOffset, Delay, LightFieldDevice + InputImageOffset);
    }
    status = cudaGetLastError();
    if (status != cudaSuccess) {
        printf("failed %s\n", cudaGetErrorString(status));
    }
    cudaDeviceSynchronize();

    if (threadIdx.x == 5)
    {
        printf("The values at beginig of %d  %d are %f\n", threadIdx.x, threadIdx.y, *(LightFieldDevice + InputImageOffset));
    }
}

子カーネル

__global__ void ConvolutionColumn(DataIn DataInput,float* __restrict__ DeviceInput, float Delay, float* __restrict__ DeviceResult)
{
    extern __shared__ float ConvolutionBlockLeft[];
    int BlockStart = blockDim.y*blockIdx.y*DataInput.V + blockIdx.x*blockDim.x;
    //int BlockEnd = BlockStart+(blockDim.x*blockDim.y)-1;
    int PixelId = blockDim.x*threadIdx.y + threadIdx.x; //32 by 24 kernal
    int LoadPixelId = DataInput.V*threadIdx.y + threadIdx.x;
    int LoadLeft,LoadRght,LoadCentre;
    float KernalSum;
    float DelayPower = Delay;
    //load upper values
    if (blockIdx.y == 0)
    {
        LoadLeft = DataInput.V*(blockDim.y - threadIdx.y-1) + threadIdx.x;
    }
    else
    {
        LoadLeft =  LoadPixelId - (DataInput.V*blockDim.y);
    }
    *(ConvolutionBlockLeft + (threadIdx.y*blockDim.x) + threadIdx.x) = *(DeviceInput + BlockStart + LoadLeft);
    if (blockIdx.y*blockDim.y + threadIdx.y >= DataInput.U)
    {
        LoadCentre = ((DataInput.U - 1)*DataInput.V) + (blockDim.x*blockIdx.x) + threadIdx.x - ((blockIdx.y*blockDim.y + threadIdx.y) - DataInput.U)*DataInput.V;
    }
    else
    {
        LoadCentre = BlockStart+LoadPixelId;
    }
    *(ConvolutionBlockLeft + (blockDim.x*blockDim.y) + (threadIdx.y*blockDim.x) + threadIdx.x) = *(DeviceInput + LoadCentre);
    if (blockIdx.y*blockDim.y + threadIdx.y + blockDim.y >= DataInput.U)
    {
        LoadRght = ((DataInput.U - 1)*DataInput.V) + (blockDim.x*blockIdx.x) + threadIdx.x - ((((blockIdx.y*blockDim.y) + threadIdx.y + blockDim.y) - DataInput.U)*DataInput.V);
    }
    else
    {
        LoadRght = BlockStart+LoadPixelId + (DataInput.V*blockDim.y);
    }
    //float tempfil, tempdata;
    //int t;
    *(ConvolutionBlockLeft + (2 * blockDim.x*blockDim.y) + (threadIdx.y*blockDim.x) + threadIdx.x) = *(DeviceInput + LoadRght);
    __syncthreads();
    float FilterSum = *(ConvolutionBlockLeft + ((blockDim.x*blockDim.y) + PixelId));
    for (int k = 1; k < DataInput.KernalNoOfFilters; k++)
    {
        KernalSum = 0;
        //printf("The value of filter size is %d\n", (DeviceFilterSize[k]));
        for (int l = -((*(DeviceFilterSize + k) - 1) / 2); l < ((*(DeviceFilterSize + k) + 1) / 2); l++)
        {
            //tempfil = *(DeviceFilterKernal + k*DataInput.KernalFilterLength + ((*(DeviceFilterSize + k) - 1) / 2) + l);
            //t = (blockDim.x*blockDim.y) + PixelId + (l*blockDim.x);
            //tempdata = *(ConvolutionBlockLeft + ((blockDim.x*blockDim.y) + PixelId - (l*blockDim.x)));
            KernalSum += *(DeviceFilterKernal + k*DataInput.KernalFilterLength + ((*(DeviceFilterSize + k) - 1) / 2) + l)**(ConvolutionBlockLeft + ((blockDim.x*blockDim.y) + PixelId - (l*blockDim.x)));
        }
        KernalSum *= DelayPower;
        DelayPower *= Delay;
        FilterSum += KernalSum;
    }
    if (blockIdx.y*blockDim.y + threadIdx.y < DataInput.U)
    {
        *(DeviceResult + LoadPixelId + BlockStart) = FilterSum;
    }
}

ここでは、子カーネルだけで問題なく動作します。ただし、別のカーネルから起動された場合、未指定の起動失敗エラーでホストから親カーネルが起動された後、cudaDeviceSynchronize()エラーが発生します (エラーはカーネル内の printf から出力されません)。

親カーネルの起動構成は<<<1,(17 17)>>>. 親からの 1 つのスレッドのみが子グリッドの起動を許可されている場合、コードは正常に機能します。1 つのブロックから起動できるグリッドの数に制限はありますか?