cuda - tex 对象访问总是返回零——有什么想法吗？

Question

我正在运行 CUDA 5.0，使用 670 设置了 compute_30、sm_30。

我通过以下方式创建一个 mipmapped 数组:

cudaExtent size;
size.width = window_width;      // 600
size.height = window_height;    // 600
size.depth = 1;
int levels = getMipMapLevels(size);
levels = MIN(levels, 9);        // 9
cudaChannelFormatDesc fp32;
fp32.f = cudaChannelFormatKindFloat;
fp32.x = fp32.y = fp32.z = fp32.w = 32;
cudaMipmappedArray_t A;
checkCuda(cudaMallocMipmappedArray(&A, &fp32, size, levels, cudaArraySurfaceLoadStore));

我用 surf2Dwrites 加载了 A 的第一层。我知道这是可行的，因为我将该数组复制到主机并将其转储到图像文件。我现在希望用 mipmap 填充 A 的其他 miplevel。通过该循环的一次迭代看起来像:

width >>= 1; width = MAX(1, width);
height >>= 1; height = MAX(1, height);

cudaArray_t from, to;
checkCuda(cudaGetMipmappedArrayLevel(&from, A, newlevel-1));
checkCuda(cudaGetMipmappedArrayLevel(&to, A, newlevel));

cudaTextureObject_t from_texture;
create_texture_object(from, true, &from_texture);
cudaSurfaceObject_t to_surface;
create_surface_object(to, &to_surface);

dim3 blocksize(16, 16, 1);
dim3 gridsize((width+blocksize.x-1)/blocksize.x,(height+blocksize.y-1)/blocksize.y, 1);
d_mipmap<<<gridsize, blocksize>>>(to_surface, from_texture, width, height);

checkCuda(cudaDeviceSynchronize());
checkCuda(cudaGetLastError());
uncreate_texture_object(&from_texture);
uncreate_surface_object(&to_surface);

已知 create_surface_object() 代码可以工作。为了以防万一，这里是 create_texture_object() 代码:

static void create_texture_object(cudaArray_t tarray, bool filter_linear, cudaTextureObject_t *tobject)
{
    assert(tarray && tobject);
    // build the resource
    cudaResourceDesc    color_res;
    memset(&color_res, 0, sizeof(cudaResourceDesc));
    color_res.resType = cudaResourceTypeArray;
    color_res.res.array.array = tarray;

    // the texture descriptor
    cudaTextureDesc     texdesc;
    memset(&texdesc, 0, sizeof(cudaTextureDesc));
    texdesc.addressMode[0] = cudaAddressModeClamp;
    texdesc.addressMode[1] = cudaAddressModeClamp;
    texdesc.addressMode[2] = cudaAddressModeClamp;
    texdesc.filterMode = filter_linear ? cudaFilterModeLinear : cudaFilterModePoint;
    texdesc.normalizedCoords = 1;

    checkCuda(cudaCreateTextureObject(tobject, &color_res, &texdesc, NULL));
}

d_mipmap 设备函数如下:

__global__ void
d_mipmap(cudaSurfaceObject_t out, cudaTextureObject_t in, int w, int h)
{
    float x = blockIdx.x * blockDim.x + threadIdx.x;
    float y = blockIdx.y * blockDim.y + threadIdx.y;

    float dx = 1.0/float(w);
    float dy = 1.0/float(h);

    if ((x < w) && (y < h))
    {
#if 0
        float4 color = 
            (tex2D<float4>(in, (x + .25f) * dx, (y + .25f) * dy)) +
            (tex2D<float4>(in, (x + .75f) * dx, (y + .25f) * dy)) +
            (tex2D<float4>(in, (x + .25f) * dx, (y + .75f) * dy)) +
            (tex2D<float4>(in, (x + .75f) * dx, (y + .75f) * dy));
        color /= 4.0f;
        surf2Dwrite(color, mipOutput, x * sizeof(float4), y);
#endif
     float4 color0 = tex2D<float4>(in, (x + .25f) * dx, (y + .25f) * dy);
     surf2Dwrite(color0, out, x * sizeof(float4), y);
    }
}

其中包含 mipmap 采样代码(如果有的话)和调试代码。

问题是，color0 始终统一为零，我一直无法理解为什么。我已经将过滤更改为点(从线性)但没有成功。我检查过错误。什么都没有。

我在这里使用 CUDA/OpenGL 互操作，但 mipmap 生成仅在 CUDA 数组上完成。

我真的不想使用纹理引用。

有什么关于去哪里看的建议吗？

Answer 1

这个错误原来是 cudaMipmappedArrays 的使用(数组或纹理对象——我无法判断哪个损坏了。)

当我修改代码以仅使用 cudaArrays 时，纹理引用再次开始工作。

由于无绑定(bind)纹理程序示例有效，该错误似乎仅限于 float32 channel mipmapped 纹理。 (我有一个测试程序显示 1 channel 和 4 channel float32 mipmapped 纹理都会出现错误。)

我已将错误报告给 Nvidia。