3

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

我通过以下方式创建了一个 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 数组上完成。

我真的真的不想使用纹理参考。

关于在哪里看的任何建议？

4

1 回答 1

4

该错误原来是使用 cudaMipmappedArrays （数组或纹理对象 - 我无法判断哪个坏了。）

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

由于无绑定纹理程序示例有效，因此该错误似乎仅限于 float32 通道 mipmaped 纹理。（我有一个测试程序，显示 1 通道和 4 通道 float32 mipmapped 纹理都会出现该错误。）

我已经向 Nvidia 报告了这个错误。

于 2013-04-05T05:07:10.820 回答