CUDA PTX f32.f32 纹理读取
Posted
技术标签:
【中文标题】CUDA PTX f32.f32 纹理读取【英文标题】:CUDA PTX f32.f32 texture read 【发布时间】:2016-01-29 10:02:27 【问题描述】:是否可以直接使用浮点索引从 CUDA 纹理中读取,例如我可以使用tex.1d.v4.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
这是可能的还是我找错了树?
【问题讨论】:
您为什么使用v4
变体?这是偶然的还是故意的?
别无选择。对于不使用 FP16 的 1d 纹理查找,纹理加载始终返回 32 位值的 4 元素向量。见这里:docs.nvidia.com/cuda/parallel-thread-execution/…
确实如此。你每天学习新的东西。我只写过用于表面访问的 PTX,其中有 .none
、.v2
和 .v4
修饰符,我只是假设标准纹理指令是相同的
【参考方案1】:
您的问题是您正在使用不受支持的纹理指令,该指令以默认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;
即。将整数索引传递给纹理单元,而不是浮点数,我想你会发现它会正常工作。您需要具有过滤读取模式的纹理才能使用tex.1d.v4.f32.f32
。
【讨论】:
对于cudaReadModeElementType
的读取模式,不进行过滤,唯一有效的坐标类型是整数索引。你会发现tex1DFetch
在这种情况下也会发出tex.1d.v4.f32.s32
。以上是关于CUDA PTX f32.f32 纹理读取的主要内容,如果未能解决你的问题,请参考以下文章