内联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) : "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”匹配哪种操作则有前面的指令决定。
也可在内联汇编里声明局部变量:
asm( “reg.u32 a;/n/t”
“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(当然,这里所说的一些细节手册并未提到).
好了,写的比较仓促,且耐心不足,疏漏之处在所难免,欢迎指正。以后会补上更详细晚上的PTX内联汇编编程文档。
欢迎大家来和我讨论交流。
下面提上完整的测试程序:
内核代码正式上面的”cuk_lerp”,但注意:测试时须将cuk_lerp放入extern “C” {}中。
另外设置编译选项时,输出不能设置为’-ptx’,只有’-cubin’或者’-fatbin’选项才支持内联ptx.
host code:
#include<stdio.h>
#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” );