CUDAのさまざまな種類のテクスチャをテクスチャ参照にバインドする方法は？

Question

コードのこの作品は、私は、テクスチャからの任意のアイデアをテクスチャの異なる種類を取得することができない文法を変更したクーダのでクーダ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 テクスチャオブジェクトを参照していますが、それはすべての発生を変更するための多くの作業です。マイナーなコード変更でより良い解決策がありますか？

おかげ

誰でも元のコードをしたい場合は、hereをクリックしてください。

Answer 1

このため最小再現ケースがあるように思える：

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): error: cannot take address of texture/surface variable "__tex0" in __device__/__global__ functions

私はこれが正しいと信じています。テクスチャ参照は、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レベルでの読み取りがテクスチャので作品余分なベクトル要素が空で無視されても、参照読み取りは常に4つのワイドベクトル型を返します。私はこれがコンパイラのバグとして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ツールチェーンと、これが発光する：放出されたCUDA 4.2ツールチェーンと同じであるPTX

// 
// 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; 
} 

を。これはコンパイラがほぼ同じレベルの型安全性チェックをインラインPTXに適用することができないためです。しかし、あなたが本当にこれをしたいかどうかについて懸命に考えてください、それは（私の意見では）未定義の行動であるからです。

テクスチャ参照がPTXで処理される方法のため、明示的な引数として渡すことはできないため、コード内のテクスチャごとに1つの読み取り関数を定義する必要があります。