如何在CUDA中将不同种类的纹理绑定到纹理引用
How to bind different kinds of textures to a texture reference in CUDA?
这段代码适用于Cuda 4.2
extern "C" texture<int,1,cudaReadModeElementType> __tex0;
extern "C" __global__ void kernel(){
float4 f = tex1Dfetch(*(texture<float4,1,cudaReadModeElementType>*)&__tex0,ii_z)
}
由于库达改变了语法,我无法从一个纹理中提取不同类型的纹理,知道吗?
PS。我在参考中找到了Cuda纹理对象,但要改变所有出现的情况需要做很多工作。有没有一个较小的代码更改的更好的解决方案?
谢谢
如果有人想要原始代码,请单击此处。
这似乎是最小的repro情况是:
texture<int,1,cudaReadModeElementType> __tex0;
__global__ void kernel0(float4 *out)
{
int t__a = blockIdx.x*blockDim.x+threadIdx.x;
int ii = (t__a*3);
float4 rr = tex1Dfetch(*(texture<float4,1,cudaReadModeElementType>*)&__tex0,ii);
out[t__a] = rr;
}
CUDA 7.5将无法编译此内核,并出现错误:
texture_repo.cu(7):错误:无法获取纹理/表面的地址
__device__/__global__
函数中的变量"__tex0"
我相信这是正确的。纹理引用是不透明的占位符类型,它没有POD类型的任何常见属性,我会非常怀疑是否会像您提供的链接那样编写代码。
然而,CUDA 4.2确实会编译它并发出有效的PTX:
.entry _Z7kernel0P6float4(
.param .u64 _Z7kernel0P6float4_param_0
)
{
.reg .f32 %f<25>;
.reg .s32 %r<8>;
.reg .s64 %rl<5>;
ld.param.u64 %rl1, [_Z7kernel0P6float4_param_0];
cvta.to.global.u64 %rl2, %rl1;
.loc 2 5 1
mov.u32 %r2, %ntid.x;
mov.u32 %r3, %ctaid.x;
mov.u32 %r4, %tid.x;
mad.lo.s32 %r5, %r2, %r3, %r4;
.loc 2 6 1
mul.lo.s32 %r1, %r5, 3;
mov.u32 %r6, 0;
// inline asm
tex.1d.v4.f32.s32 {%f1, %f2, %f3, %f4}, [__tex0, {%r1}];
// inline asm
.loc 2 8 1
mul.wide.s32 %rl3, %r5, 16;
add.s64 %rl4, %rl2, %rl3;
st.global.v4.f32 [%rl4], {%f1, %f2, %f3, %f4};
.loc 2 9 2
ret;
}
强制转换显然除了抑制编译器错误之外没有任何效果,而且在PTX级别,读取是有效的,因为纹理引用读取总是返回一个四宽的向量类型,即使额外的向量元素是空的并被忽略。我认为这是CUDA 4.2中编译的一个编译器错误,在这种情况下,CUDA 7.5似乎是正确的。
也就是说,一个非常棘手的解决方案是这样做:
texture<int,1,cudaReadModeElementType> __tex0;
__device__ float4 tex_load0(int idx)
{
float4 temp;
asm("tex.1d.v4.f32.s32 {%0, %1, %2, %3}, [__tex0, {%4}];" :
"=f"(temp.x), "=f"(temp.y), "=f"(temp.z), "=f"(temp.w) : "r"(idx));
return temp;
}
__global__ void kernel1(float4 *out)
{
int t__a = blockIdx.x*blockDim.x+threadIdx.x;
int ii = (t__a*3);
float4 rr = tex_load0(ii);
out[t__a] = rr;
}
[免责声明:已编译但从未测试过。不建议使用。使用风险自负]。
即将CUDA 4.2编译器内联发出的相同PTX插入到设备函数中,并用对设备函数的调用替换纹理获取。对于CUDA 7.5工具链,这会发出:
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-19856038
// Cuda compilation tools, release 7.5, V7.5.17
// Based on LLVM 3.4svn
//
.version 4.3
.target sm_30
.address_size 64
// .globl _Z9tex_load0i
.global .texref __tex0;
.visible .func (.param .align 16 .b8 func_retval0[16]) _Z9tex_load0i(
.param .b32 _Z9tex_load0i_param_0
)
{
.reg .f32 %f<5>;
.reg .b32 %r<2>;
ld.param.u32 %r1, [_Z9tex_load0i_param_0];
// inline asm
tex.1d.v4.f32.s32 {%f1, %f2, %f3, %f4}, [__tex0, {%r1}];
// inline asm
st.param.f32 [func_retval0+0], %f1;
st.param.f32 [func_retval0+4], %f2;
st.param.f32 [func_retval0+8], %f3;
st.param.f32 [func_retval0+12], %f4;
ret;
}
// .globl _Z7kernel1P6float4
.visible .entry _Z7kernel1P6float4(
.param .u64 _Z7kernel1P6float4_param_0
)
{
.reg .f32 %f<5>;
.reg .b32 %r<6>;
.reg .b64 %rd<5>;
ld.param.u64 %rd1, [_Z7kernel1P6float4_param_0];
cvta.to.global.u64 %rd2, %rd1;
mov.u32 %r2, %ctaid.x;
mov.u32 %r3, %ntid.x;
mov.u32 %r4, %tid.x;
mad.lo.s32 %r5, %r3, %r2, %r4;
mul.lo.s32 %r1, %r5, 3;
mul.wide.s32 %rd3, %r5, 16;
add.s64 %rd4, %rd2, %rd3;
// inline asm
tex.1d.v4.f32.s32 {%f1, %f2, %f3, %f4}, [__tex0, {%r1}];
// inline asm
st.global.v4.f32 [%rd4], {%f1, %f2, %f3, %f4};
ret;
}
这与发出的CUDA 4.2工具链是相同的PTX。这是因为编译器不能将几乎相同级别的类型安全检查应用于内联PTX。但仔细想想你是否真的想这样做,因为(在我看来)这是一种未定义的行为。
还要注意的是,由于在PTX中处理纹理引用的方式,您不能将它们作为显式参数传递,因此您需要在代码中为每个纹理定义一个读取函数。
- 更正GLSL无绑定纹理句柄中的结构布局
- OpenGL - 绑定到默认纹理的做法?
- 将多个纹理 OpenGL 绑定到不同的四边形
- 为每个对象绑定不同的纹理
- 纹理绑定不起作用 / C++ / OpenGL
- OpenGL 加载将上次加载的纹理绑定到所有 textureID
- OpenGL-重新绑定已经绑定的纹理
- 在C++,OpenGL,Glut,如何将image.c绑定到纹理,其中image.c来自Gimp>Export>C源代码
- 如何绑定纹理openGL的一部分
- OpenGL 纹理绑定到对象
- 将QImage绑定到纹理以在QOpenGLWidget中显示视频的正确方法是什么
- 将绑定到FBO的纹理复制到另一个OpenGL上下文
- Qt 5 OpenGL绑定纹理导致奇怪的错误
- 纹理绑定:OpenGL+OpenCV
- 在OpenGL中绑定多于MAX_TEXTURE_UNITS的纹理
- QOpenGLFramebufferObject绑定纹理
- 在重复模式下绑定纹理
- OpenGL绑定几个2D纹理
- 如何在CUDA中将不同种类的纹理绑定到纹理引用
- OpenGL 将纹理绑定到顶点缓冲区对象(使用 CG 着色器)