Texture objects fetching with unknown texture address

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi!
I’m trying to understand how are texture objects instructions translated into CUDA assembly code. If we use different textures in different threads, than generated code includes awful strange instructions (even with synchronization).

Here is simple example:

$ cat test.cu
__global__ void kernel(float* a, const cudaTextureObject_t* tex){
  a[0] = tex3D<float>(tex[threadIdx.x], 0.1, 0.2, 0.3);
}

It translates to this(for arch=sm_50):

.text._Z6kernelPfPKy:
        /*0008*/                   MOV R1, c[0x0][0x20];
        /*0010*/                   S2R R0, SR_TID.X;
        /*0018*/                   MOV32I R4, 0x3dcccccd;
        /*0028*/                   MOV32I R5, 0x3e4ccccd;
        /*0030*/                   MOV32I R6, 0x3e99999a;
        /*0038*/                   SHL.W R2, R0.reuse, 0x3;
        /*0048*/                   SHR.U32 R0, R0, 0x1d;
        /*0050*/                   IADD R2.CC, R2, c[0x0][0x148];
        /*0058*/                   IADD.X R3, R0, c[0x0][0x14c];
        /*0068*/                   LDG.E R0, [R2];
        /*0070*/                   DEPBAR {5,0};
        /*0078*/                   SHFL.IDX PT, R2, R0, 0x0, 0x1c1f;
        /*0088*/                   SSY `(.L_1);
        /*0090*/                   IADD RZ.CC, R2, -R0;
        /*0098*/                   BRA CC.EQ, `(.L_2);
        /*00a8*/                   SHFL.IDX PT, R2, R0, 0x1, 0x1c1f;
        /*00b0*/                   IADD RZ.CC, R2, -R0;
        /*00b8*/                   BRA CC.EQ, `(.L_2);
        /*00c8*/                   SHFL.IDX PT, R2, R0, 0x2, 0x1c1f;
        /*00d0*/                   IADD RZ.CC, R2, -R0;
        /*00d8*/                   BRA CC.EQ, `(.L_2);
.L_2:
        /*00e8*/                   TEX.B.P R0, R4, R0, 0x0, 3D, 0x1;
        /*00f0*/                   SYNC                             (*"TARGET= .L_1 "*);
.L_1:
        /*00f8*/                   MOV R2, c[0x0][0x140];
        /*0108*/                   MOV R3, c[0x0][0x144];
        /*0110*/                   STG.E [R2], R0;
        /*0118*/                   EXIT;

Many freaky instructions (0x0070-0x00f0) occur for arch=50 and lower, but for arch=52 it looks qute fine, tex fetching translates into one command: TEX.B.NODEP.P R0, R4, R0, 0x0, 3D, 0x1;