【IT168 技术】内联PTX汇编具有如下形式：
asm(“instop”::”type_symbolic”(r));
第二种形式是无输出操作格式，需要使用”::”指示符
其中instop是指令操作
type_symbolic是类型指示符(可选如下)，分别对应与PTX中的数据类型：
“h” .s16, .u16
“r” .s32, .u32
“l” .s64, .u64
“f” .f32
“d” .f64
例如：
float a=…
float b=…;
float c;
asm( “add.f32 %0, %1, %2;” : “=f”(c):”f”(a),”f”(b));
%0, %1, %2,是匹配符，在分开写的“asm()”段中，不通指令序列中的%匹配符不具有相关项，它们的作用只是根据“：”后面的匹配格式按照顺序进行匹配，所以统一规格程序中的两段“asm()“中的相同的%numberic不一定指向统一规格实际的物理寄存器，例如：
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) : "f"(a) ); // u*=a
asm( "fma.rn.f32 %0, %1, %2, %3;" : "=f"(b) : "f"(alpha), "f"(v), "f"(u) ); // b=alpha*v+u
z[ threadIdx.x ]=b;
}
来看下这段代码，首先看第一段“asm()”,在前面的指令序列中可以直接使用数字，但有些限制，0f3f800000对应的十进制浮点数是1.f,但不能直接使用”1.f”,否则编译器会报错，因为’f’属于类型匹配符中的”关键字”;也不能使用1，或者1.0，这样编译器也会报错，前者认为是整数，类型不匹配;后者则认为是双精度浮点数，类型的尺寸不匹配。
再看第二段”asm()”,“+f”(u)表示即读友写操作，并以一定对应“+=”操作，也可是任何CUDA编译器支持的“op=”操作，比如：+=, -=, *=, &=, ”op”匹配哪种操作则有前面的指令决定。
也可在内联汇编里声明局部变量：
“shl.u32 %0, 1, a;”
: “=r”(mask) : “r”(a) );
注意，这段代码，当指令操作位于匹配格式序列之前也就是最后一段指令操作学列时，不需要再使用“/n/t”换行符。
匹配格式也支持如果没有输入的操作:
asm( “mov.s32 %0, 7;” : “=r”(x) );
通常存储器写操作是作为输出操作，但有时会存在同步隐患，或者想避免编译器对存储操作的优化，这时可以使用”memory”指示字：
总体来说inline PTX现在还比较初级，有些功能还不能使用，比如指令操作数只能是标量，不支持矢量，举个例子：
asm( "ld.shared.v2.f32 { %0, %1 },[ %2+16 ];":"=f"(a),"=f"(b): “r”(ptr) );
这样虽然编译可以通过，但是内核执行却会发生错误，而应该使用如下代码代替：
asm( "ld.shared.f32 { %0 },[ %1+0 ];":"=f"(a): “r”(ptr) );
asm( "ld.shared.f32 { %0 },[ %1+8 ];":"=f"(b): “r”(ptr) );
关于使用inline ptx的更多细节可以参考CUDAtookit4.0中的using inline PTX assembly in CUDA.pdf(当然，这里所说的一些细节手册并未提到).
下面提上完整的测试程序：
内核代码正式上面的”cuk_lerp”,但注意：测试时须将cuk_lerp放入extern “C” {}中。
另外设置编译选项时，输出不能设置为’-ptx’,只有’-cubin’或者’-fatbin’选项才支持内联ptx.
host code:
#include<cuda.h>
#pragma comment( lib, "cuda.lib" )
int main()
{
CUdevice device;
CUcontext context;
CUmodule module;
CUfunction kernel;
CUdeviceptr dptr[ 3 ];
cuInit( 0 );
cuDeviceGet( &device, 0 );
cuCtxCreate( &context, CU_CTX_SCHED_AUTO, device );
cuModuleLoad( &module, "kernel.cubin" );
cuModuleGetFunction( &kernel, module, "cuk_lerp" );
#define n_threads 128
size_t size=n_threads*sizeof( float );
cuMemAlloc( &dptr[ 0 ], size );
cuMemAlloc( &dptr[ 1 ], size );
cuMemAlloc( &dptr[ 2 ], size );
float a[ 128 ];
float b[ 128 ];
for( int i=0; i<n_threads; ++i )
{
a[ i ]=1.f;
b[ i ]=2.f;
}
cuMemcpyHtoD( dptr[ 1 ], a, size );
cuMemcpyHtoD( dptr[ 2 ], b, size );
float alpha=0.5f;
void* params[]={ &dptr[ 0 ], &dptr[ 1 ], &dptr[ 2 ], &alpha };
cuLaunchKernel( kernel, 1, 1, 1, 128, 1, 1, 0, NULL, params, 0 );
cuCtxSynchronize();
cuMemcpyDtoH( a, dptr[ 0 ], size );
for( int i=0; i<128; ++i ){
printf( "%f/n", a[ i ] );
}
cuMemFree( dptr[ 0 ] );
cuMemFree( dptr[ 1 ] );
cuMemFree( dptr[ 2 ] );
cuModuleUnload( module );
cuCtxDestroy( context );
return 0;
}
asm volatile ( “mov.u32 %0, %%laneid.x;” : “=r”(out) :: “memory” );
更多内容请点击：
CUDA专区：http://cuda.it168.com/
CUDA论坛：http://cudabbs.it168.com/
