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

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

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

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

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

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

система: GTX Titan Black (см_3,5), CUDA6.0.

РЕДАКТИРОВАТЬ: пример кода для иллюстрации.

Очевидно, что Efield объявлен и заполнен ранее. HANDLE_ERROR происходит из книги. 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;
}