Wie binden Sie verschiedene Arten von Texturen an eine Texturreferenz in CUDA?

Question

Dieses Stück Code funktioniert auf 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) 
} 

Seit Cuda grammer haben sich geändert, ich nicht andere Art von Texturen aus einer Textur holen kann, eine Idee?

PS. Ich habe Cuda Textur Objekt in Bezug gefunden, aber das ist eine Menge Arbeit, alle Vorkommen zu ändern. Gibt es eine bessere Lösung mit geringfügiger Codeänderung?

Dank

Wenn jemand den ursprünglichen Code möchten, benutzen Sie bitte here klicken.

Answer 1

Es scheint, wie der minimale repro Fall hierfür ist:

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 fehl diesen Kernel mit einem Fehler zu kompilieren:

texture_repo.cu(7): error: cannot take address of texture/surface variable "__tex0" in __device__/__global__ functions

Ich glaube, das ist richtig. Textur-Referenzen sind undurchsichtige Platzhalter-Typen, die alle üblichen Eigenschaften von POD-Typen nicht haben und ich würde sehr verdächtig sein über jemals Code wie das Beispiel schreiben Sie einen Link zu versehen.

Es ist jedoch wahr, dass CUDA 4.2 diese kompiliert und emittieren gültig 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; 
} 

Die Besetzung hat anscheinend keinen anderen Effekt als einen Compiler-Fehler zu unterdrücken, und auf einem Niveau PTX die Lese funktioniert, weil Textur Referenzlesevorgänge geben immer einen vier breiten Vektortyp zurück, auch wenn die zusätzlichen Vektorelemente leer sind und ignoriert werden. Ich würde die Tatsache betrachten, dass dies in CUDA 4.2 als Compiler-Fehler kompiliert wird, und es scheint, dass CUDA 7.5 in diesem Fall korrekt ist.

Das heißt, ein sehr hacky Behelfslösung wäre, dies zu tun:

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

[HAFTUNGSAUSSCHLUSS: kompiliert aber nie getestet. Nicht empfohlen. Verwendung auf eigene Gefahr].

d.h. einzufügen denselben PTX emittiert inline durch die CUDA 4.2 Compiler in eine Gerätefunktion, und ersetzen die Textur mit Aufrufen an die Gerätefunktion abruft. Mit der CUDA 7.5 Werkzeugkette, diese aussendet:

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

, die gleich wie die PTX CUDA 4.2 Werkzeugkette emittiert wird. Dies funktioniert, weil der Compiler fast nicht die gleiche Sicherheitsstufe für den Typ auf Inline-PTX anwenden kann. Aber denken Sie darüber nach, ob Sie das wirklich wollen, weil es (meiner Meinung nach) undefiniertes Verhalten ist.

Beachten Sie auch, dass wegen der Art, Beschaffenheit Referenzen in PTX behandelt werden, können Sie sie nicht als explizite Argumente übergeben können, so dass Sie benötigen in Ihrem Code eine Lesefunktion pro Textur definieren.