CUB的TexRefInputIterator是如何工作的?
How does CUB's TexRefInputIterator work?
CUB为纹理引用提供了一个迭代器,其实现很容易访问。
因为我不能弄清楚如何实现可模板的纹理引用自己-他们"只能声明为一个静态全局变量" -我现在试图理解它是如何在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 %dn", 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()));
}
我得到的最接近于下面的东西,基于通常如何访问静态模板成员。为了清楚起见,下面简单地从上面删除了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 %dn", 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
该编译错误是由于使用包含模板类型的宏生成的代码,因此模板中的逗号使预处理器认为它们是更多的参数。我通过修补crt/host_runtime头并使这些宏(__text_var, __device__text_var和__name__text_var)的cpp参数变进来解决这个问题。换句话说,将cpp替换为cpp....
相关文章:
- QSqlquery prepare()和bindvalue()不工作
- 导入库可以跨dll版本工作吗
- 以螺旋方式打印矩阵的程序.(工作不好)
- 对象指针在c++中是如何工作的
- 为什么在Windows上的VS 2019和Clang 9中"size_t"在没有标题的情况下工作
- VSOMEIP-2个设备之间的通信(TCP/UDP)不工作
- 为字符串中每 N 个字符插入空格的函数没有按照我认为的方式工作?
- C++为线程工作动态地分割例程
- 为什么我的 std::ref 无法按预期工作?
- 布尔比较运算符是如何在C++中工作的
- SampleConsensusPrerejective(ext.RANSAC)是如何真正工作的
- 不确定要在我的main中放入什么才能使我的代码正常工作
- 为什么std::condition_variable notify_all的工作速度比notify_one快(对于随机请
- <<操作员在下面的行中工作
- 有人能解释一下为什么下界是这样工作的吗C++的
- ExtractIconEx:可以工作,但偶尔会崩溃
- C++中的memset函数工作不正常
- 当我在第一个循环中使用"auto"时,它工作正常,但是使用"int"它会给出错误,为什么?
- 当 int 方法工作正常时,void 方法有何不同,或者为什么我不能调用 void 方法?
- sdl软件渲染器不工作,工作在硬件加速的一个