CUDA4.0 inline PTX汇编程序开发( 0 )

内联PTX汇编具有如下形式：

asm(“instop”:”type_symbolic”(or):”type_symbolic”(ir),..);

asm(“instop”::”type_symbolic”(r));

第二种形式是无输出操作格式，需要使用”::”指示符

其中instop是指令操作

type_symbolic是类型指示符（可选如下），分别对应与PTX中的数据类型：

“h” .s16, .u16

“r” .s32, .u32

“l” .s64, .u64

“f” .f32

“d” .f64

 

例如：

//c=a+b

float a=…

float b=…;

float c;

asm( “add.f32 %0, %1, %2;” : “=f”(c):”f”(a),”f”(b));

 

%0, %1, %2,是匹配符，在分开写的“asm（）”段中，不通指令序列中的%匹配符不具有相关项，它们的作用只是根据“：”后面的匹配格式按照顺序进行匹配，所以统一规格程序中的两段“asm()“中的相同的%numberic不一定指向统一规格实际的物理寄存器，例如：

 

__global__

void cuk_lerp( float* z, const float* x, const float* y, float alpha )

{

    float u=x[ threadIdx.x ];

    float v=y[ threadIdx.x ];

    float a, b;

    asm( "sub.f32 %0, 0f3f800000, %1;" : "=f"(a) : "f"(alpha) );  //a=1.f-alpha

    asm( "mul.f32 %0, %0, %1;" : "+f"(u) :