CUDA4.0 inline PTX汇编程序开发

内联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” );