CUDA PTX f32.f32纹理读取

是否可以直接使用浮点索引从CUDA纹理读取数据，例如：我可以使用tex.1d.v4.f32.f32来执行纹理获取。CUDA PTX f32.f32纹理读取

在查看.ptx文件时，这似乎节省了两条指令，这反映在进行基准测试时性能提高。然而，相当关键的缺点是，虽然这看起来没有问题，但它不会产生预期的结果。

下面的代码演示了这个问题：

#include "cuda.h" 
#include <thrust/device_vector.h> 

//create a global 1D texture of type float 
texture<float, cudaTextureType1D, cudaReadModeElementType> tex; 

//below is a hand rolled ptx texture lookup using tex.1d.v4.f32.f32 
__device__ 
float tex_load(float idx) 
{ 
    float4 temp; 
    asm("tex.1d.v4.f32.f32 {%0, %1, %2, %3}, [tex, {%4}];" : 
     "=f"(temp.x), "=f"(temp.y), "=f"(temp.z), "=f"(temp.w) : "f"(idx)); 
    return temp.x; 
} 

//Try to read from the texture using tex1Dfetch and the custom tex_load 
__global__ void read(){ 
    float x = tex1Dfetch(tex,0.0f); 
    float y = tex_load(0.0f); 
    printf("tex1Dfetch: %f tex_load: %f\n",x,y); 
} 

int main() 
{ 
    //create a vector of size 1 with the x[0]=3.14 
    thrust::device_vector<float> x(1,3.14); 
    float* x_ptr = thrust::raw_pointer_cast(&x[0]); 

    //bind the texture 
    cudaBindTexture(0, tex, x_ptr, sizeof(float)); 

    //launch a single thread single block kernel 
    read<<<1,1>>>(); 
    cudaUnbindTexture(tex); 
    return 0; 
}

我已经试过这对情侣卡（K40，C2070），并与一对夫妇CUDA的版本（6.0,7.0），但所有我得到相同的输出：

tex1Dfetch: 3.140000 tex_load: 0.000000

这是可能的还是我吠叫错了树？

来源

2016-01-29 ebarr

你为什么使用'v4'变种？那是偶然的还是故意的？ – talonmies

没有其他选择。对于不使用FP16的1d纹理查找，纹理加载总是返回32位值的4元素向量。看到这里：http://docs.nvidia.com/cuda/parallel-thread-execution/#texture-instructions-tex – ebarr

确实。你每天学习新的东西。我只写过用于表面访问的PTX，其中有'.none'，'.v2'和'.v4'修饰符，我只是假设标准纹理指令将是相同的 – talonmies

你的问题是，你正在使用一个不受支持的指令来使用默认的cudaReadModeElementType读模式绑定到线性内存的纹理。如果你改写你的功能：

__device__ 
float tex_load(int idx) 
{ 
    float4 temp; 
    asm("tex.1d.v4.f32.s32 {%0, %1, %2, %3}, [tex, {%4}];" : 
     "=f"(temp.x), "=f"(temp.y), "=f"(temp.z), "=f"(temp.w) : "r"(idx)); 
    return temp.x; 
}

ie。传递一个整数索引到纹理单元，而不是一个浮点数，我想你会发现它会正常工作。您需要使用过滤读取模式的纹理，以使用tex.1d.v4.f32.f32。

来源

2016-01-31 11:35:10 talonmies

对于读取模式'cudaReadModeElementType' ，不执行过滤，唯一有效的坐标类型是整数索引。在这种情况下，你会发现'tex1DFetch'将会发出'tex.1d.v4.f32.s32'。 – talonmies

CUDA PTX f32.f32纹理读取

回答

相关问题