c++ - CUB 的 TexRefInputIterator 是如何工作的？

Question

CUB provides an iterator对于纹理引用，其中的实现 is readily accessible .

因为我不知道如何自己实现可模板化的纹理引用 - 他们 "can only be declared as a static global variable" - 我现在正试图了解它是如何在 CUB 中完成的。但其中一些超出了我的 C++ 知识范围，而且我无法在其他地方找到答案(话又说回来，我真的不知道要搜索什么)。

具体来说:

IteratorTexRef 周围未命名的namespace 有意义吗？ 我只能认为是限制IteratorTexRef::TexId::ref 到文件/翻译单元范围。

IteratorTexRef 的用途是什么？ 它只包装 TexId，但删除它会导致(对我而言)难以理解的编译时错误.

此代码是链接到实现的精简版本，可编译并运行:

#include <thrust/device_vector.h>

namespace {

template <typename T>
struct IteratorTexRef
{
    template <int UNIQUE_ID>
    struct TexId
    {
        // Assume T is a valid texture word size.
        typedef texture<T> TexRef;

        static TexRef ref;

        static __device__ T fetch(ptrdiff_t offset)
        {
            return tex1Dfetch(ref, offset);
        }
    };
};

template <typename  T>
template <int UNIQUE_ID>
typename IteratorTexRef<T>:: template TexId<UNIQUE_ID>::TexRef IteratorTexRef<T>:: template TexId<UNIQUE_ID>::ref;

} // Anomymous namespace

template <typename T, int UNIQUE_ID = 0>
class TextureRefIterator
{
private:
    typedef typename IteratorTexRef<T>:: template TexId<UNIQUE_ID> TexId;
    ptrdiff_t tex_offset;

public:
    __device__ T operator[](int i) const
    {
        return TexId::fetch(this->tex_offset + i);
    }

    cudaError_t bind(
        const T* const ptr,
        size_t bytes = size_t(-1))
    {
        size_t offset;
        cudaError_t state = cudaBindTexture(&offset, TexId::ref, ptr, bytes);
        this->tex_offset = (ptrdiff_t) (offset / sizeof(T));
        return state;
    }
};

template <typename TexIter>
__global__ void kernel(TexIter iter)
{
    int a = iter[threadIdx.x];
    printf("tid %d, a %d\n", threadIdx.x, a);
}

template <typename T>
void launch_kernel(T* d_in)
{
    TextureRefIterator<T> tex_iter;
    tex_iter.bind(d_in);

    kernel<<<1, 32>>>(tex_iter);
}

int main()
{
    thrust::device_vector<float> d_in(32, 1);
    launch_kernel(thrust::raw_pointer_cast(d_in.data()));
}

我得到的最接近的是类似于下面的东西，基于一个人会如何normally access a static template member .为清楚起见，下面简单地从上面删除 IteratorTexRef:

#include <thrust/device_vector.h>

namespace {

template <typename T, int UNIQUE_ID>
struct TexId
{
    // Assume T is a valid texture word size.
    typedef texture<T> TexRef;

    static TexRef ref;

    static __device__ T fetch(ptrdiff_t offset)
    {
        return tex1Dfetch(ref, offset);
    }
};

template <typename  T, int UNIQUE_ID>
typename TexId<T, UNIQUE_ID>::TexRef TexId<T, UNIQUE_ID>::ref;


} // Anonymous namespace

template <typename T, int UNIQUE_ID = 0>
class TextureRefIterator
{
private:
    typedef TexId<T, UNIQUE_ID> TexId;
    ptrdiff_t tex_offset;

public:
    __device__ T operator[](int i) const
    {
        return TexId::fetch(this->tex_offset + i);
    }

    cudaError_t bind(
        const T* const ptr,
        size_t bytes = size_t(-1))
    {
        size_t offset;
        cudaError_t state = cudaBindTexture(&offset, TexId::ref, ptr, bytes);
        this->tex_offset = (ptrdiff_t) (offset / sizeof(T));
        return state;
    }
};

template <typename TexIter>
__global__ void kernel(TexIter iter)
{
    int a = iter[0];
    printf("tid %d, a %d\n", threadIdx.x, a);
}

template <typename T>
void launch_kernel(T* d_in)
{
    TextureRefIterator<T> tex_iter;
    tex_iter.bind(d_in);

    kernel<<<1, 32>>>(tex_iter);
}

int main()
{
    thrust::device_vector<float> d_in(32, 1);
    launch_kernel(thrust::raw_pointer_cast(d_in.data()));
}

它给出了这些有些深奥的编译时错误。 (使用 nvcc iter.cu 和 CUDA 7.0 编译):

In file included from tmpxft_000057d4_00000000-4_test2.cudafe1.stub.c:1:0:
/tmp/tmpxft_000057d4_00000000-4_test2.cudafe1.stub.c:30:3737: error: macro "__text_var" passed 3 arguments, but takes just 2
 dIfLi0EE3refE,::_NV_ANON_NAMESPACE::TexId<float, (int)0> ::ref), 1, 0, 0);__cudaReg
                                                                         ^
/tmp/tmpxft_000057d4_00000000-4_test2.cudafe1.stub.c:30:1: error: macro "__device__text_var" passed 3 arguments, but takes just 2
 static void __nv_cudaEntityRegisterCallback(void **__T2202){__nv_dummy_param_ref(__
 ^
/tmp/tmpxft_000057d4_00000000-4_test2.cudafe1.stub.c:30:1: error: macro "__name__text_var" passed 3 arguments, but takes just 2

Answer 1

该编译错误是由于生成的代码使用了包含模板类型的宏，因此模板中的逗号使预处理器认为它们是更多参数。我通过修补 crt/host_runtime header 并使这些宏(__text_var、__device__text_var 和 __name__text_var)的 cpp 参数可变来解决此问题。换句话说，将 cpp 替换为 cpp....