Как работает TexRefInputIterator CUB?

Question

CUB provides an iterator для текстурных ссылок, реализация которых is readily accessible.

Поскольку я не мог понять, как реализовать шаблонные ссылки на тексту, я их "can only be declared as a static global variable" - Теперь я пытаюсь понять, как это делается в CUB. Но некоторые из них не соответствуют моим знаниям на С ++, и я не смог найти ответы в другом месте (опять же, я действительно не знаю, что искать).

В частности:

ли неназванный namespace окружающих IteratorTexRef значимым? Я могу только думать, что это ограничение 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 и сделав параметр cpp этих макросов (__text_var, __device__text_var и __name__text_var) и переменный. Другими словами, замените cpp на cpp ....