c++ - CUDA动态并行性: invalid global write when using texture memory

Question

已关闭。此问题需要 debugging details 。目前不接受答案。

---

编辑问题以包含 desired behavior, a specific problem or error, and the shortest code necessary to reproduce the problem 。这将有助于其他人回答问题。

已关闭 8 年前。

Improve this question

当内核中的内核调用(甚至递归调用)使用纹理内存来获取值时，我似乎遇到了麻烦。

如果子内核(例如不同的内核)不使用纹理内存，则一切都很好。如果我不在内核中调用内核，结果就是预期的结果。只要我使用纹理内存(在我的情况下，由于空间局部性和快速过滤而非常有用)，cuda-memcheck 就会返回“大小为 4 的 __global__ 写入无效”。

我在编程指南中看到动态并行中，使用纹理内存时必须小心，可能会导致数据不一致，但这里子内核甚至没有启动。

我尝试过将 __syncthreads() 和 cudaDeviceSynchronize 放置在调用纹理内存之前或之后，但没有任何效果。

是否有一些已经报告的案例，我是否做错了什么，或者只是您不能以这种方式使用纹理内存？

系统:gtx titan black(sm_3.5)，CUDA6.0。

编辑:一些示例代码来说明。

显然，EField是之前声明和填充的。 HANDLE_ERROR 来自 book.h 包含来自 CUDA by examples

这是一个可编译的代码:

#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;
}

Answer 1

将来，请提供完整、可编译的代码。 SO预计this 。作为不确定性的一个例子，您的内核定义是 test_cdp。从主机代码调用的内核是 test2_cdp。请不要让别人猜测你的意图，或者玩20个问题来澄清你的代码。发布一个完整、可编译的代码，无需添加或更改，即可演示该问题。这就是您的问题票数接近的原因。

我发现有 2 个问题。

如果您要解决上述问题，所编写的代码可能会导致启动无限的子内核链。看来您可能认为 oneChild 变量在父内核和子内核之间以某种方式共享。它不是。因此，每个启动的子内核都会看到 oneChild 为零，并且它将启动自己的子内核。我不知道这个序列会在哪里结束，但这不是 CDP 的明智使用。

CDP does not support来自设备启动的内核的模块范围纹理引用。使用texture objects相反。