這似乎是最小的攝製情況下是這樣的:
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水平讀的作品,因爲質地引用讀取總是返回一個四寬的向量類型,即使額外的向量元素是空的並被忽略。我認爲這個在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工具鏈,此發射:
//
// 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作爲發射的CUDA 4.2工具鏈。這是可行的,因爲編譯器不能應用幾乎相同級別的類型安全檢查來內聯PTX。但想想你是否真的想這樣做,因爲它(在我看來)是未定義的行爲。
另請注意,由於紋理引用在PTX中處理的方式,因此無法將它們作爲顯式參數傳遞,因此您需要在代碼中爲每個紋理定義一個讀取函數。
您是否有完整的例子來說明該代碼的功能?因爲我高度懷疑它出現在你發佈它的時候,原因有很多。 – talonmies
這是來自大型項目的代碼行,它絕對有效。哪一部分困擾你?@talonmies,我想添加更多的線。 – hamwj1991
幾乎所有的東西 - 紋理上的外部聲明,將紋理轉換爲不同類型,將float4紋理加載到浮動中。你的問題基本上是「這曾經工作,現在不行,我該如何解決它?」。要回答這個問題,需要一個可以編譯和運行的實際repro案例。 – talonmies