Динамический параллелизм CUDA: недопустимая глобальная запись при использовании памяти текстур

Кажется, у меня проблемы, когда вызов ядра внутри ядра (даже рекурсивный вызов) использует текстурную память для получения значения.

Если дочернее ядро, скажем другое, не использует память текстур, все в порядке. Если я не вызываю ядро в ядре, результаты будут ожидаемыми. Пока я использую текстурную память, которая в моем случае очень полезна из-за пространственной локальности и быстрой фильтрации, cuda-memcheck возвращает "Недопустимый __global__ write size 4".

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

Я пробовал __syncthreads() и cudaDeviceSynchronize, размещенные до или после вызова текстурной памяти, но ничего.

Есть ли какие-то уже сообщенные случаи, я делаю что-то неправильно или просто вы не можете использовать текстурную память таким образом?

system: gtx titan black (sm_3.5), CUDA6.0.

EDIT: некоторый пример кода для иллюстрации.

Очевидно, EField объявляется и заполняется раньше. HANDLE_ERROR поступает из книги book.h из CUDA примерами

Вот компилируемый код:

#include "cuda.h"
#include "/common/book.h"

#define DIM 2048

texture<float4, 2, cudaReadModeElementType> texEField;

__device__ int oneChild = 0;


__global__ void test_cdp( float x0, float y0 ){
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int idx = x + y * blockDim.x * gridDim.x;

    printf("Propa started from thread %d\n", idx);      
    float4 E = tex2D( texEField, x0, y0 );

    printf("E field %f -- %f\n", E.z, E.w);     
    if( oneChild < 1 ){
        test_cdp<<<1, 1>>>(x0, y0);
        oneChild++;
    }
}

int main( void ){   

    //Start of texture allocation

    float4 *EField = new float4 [DIM*DIM];
    for( int u = 0; u < DIM*DIM; u++ ){
        EField[u].x = 1.0f;
        EField[u].y = 1.0f;
        EField[u].z = 1.0f;
        EField[u].w = 1.0f;
    }   


    cudaChannelFormatDesc desc = cudaCreateChannelDesc<float4>();

    float4 *dev_EField;
    HANDLE_ERROR( cudaMalloc( (void**)&dev_EField, DIM * DIM * sizeof(float4) ) );

    HANDLE_ERROR( cudaMemcpy( dev_EField, EField, DIM * DIM * sizeof(float4), cudaMemcpyHostToDevice ) );

    HANDLE_ERROR( cudaBindTexture2D( NULL, texEField, dev_EField, desc, DIM, DIM, sizeof(float4) * DIM ) );

    texEField.addressMode[0] = cudaAddressModeWrap;
    texEField.addressMode[1] = cudaAddressModeWrap;
    texEField.filterMode = cudaFilterModeLinear;
    texEField.normalized = true;

    test_cdp<<<1, 1>>>(0.5, 0.5);

    HANDLE_ERROR( cudaFree( dev_EField ) );
    HANDLE_ERROR( cudaUnbindTexture( texEField ) );
    return 0;
}