0

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

如果子内核，比如说另一个内核，不使用纹理内存，一切都很好。如果我不在内核中调用内核，则结果是预期的。 只要我使用由于空间局部性和快速过滤而在我的情况下非常有用的纹理内存，cuda-memcheck 就会返回"Invalid __global__ write of size 4"。

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

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

是否有一些已经报告的案例，是我做错了什么还是只是你不能那样使用纹理内存？

系统：gtx泰坦黑（sm_3.5），CUDA6.0。

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

显然，之前已声明并填充了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;
}
4

1 回答 1

0

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

我可以看到 2 个问题。

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

  2. CDP不支持来自设备启动内核的模块范围纹理引用。请改用纹理对象。

于 2014-08-20T21:23:14.287 回答