內聯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__
voidcuk_lerp(float
*z,constfloat
*x,constfloat
*y,floatalpha)
來看下這段**，首先看第一段「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 ,[ %2+16 ];":"=f"(a),"=f"(b): 「r」(ptr) );
這樣雖然編譯可以通過，但是核心執行卻會發生錯誤，而應該使用如下**代替：
asm( "ld.shared.f32 ,[ %1+0 ];":"=f"(a): 「r」(ptr) );
asm( "ld.shared.f32 ,[ %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
#include
#pragmacomment(lib,"cuda.lib")
intmain()
cumemcpyhtod
(dptr[1],a,size);
cumemcpyhtod
(dptr[2],b,size);
floatalpha=
0.5f;
void
*params=;
culaunchkernel(kernel,1,1,1,128,1,1,0,null,params,0);
cuctxsynchronize
();
cumemcpydtoh
(a,dptr[0],size);
for
(inti=
0;i<
128;++i )
cumemfree
(dptr[0]);
cumemfree
(dptr[1]);
cumemfree
(dptr[2]);
cumoduleunload
(module);
cuctxdestroy
(context);
return0; }
asm
volatile ( 「mov.u32 %0, %%laneid.x;」 : 「=r」(out) :: 「memory」 );
匯程式設計序 退出
作為第乙個匯程式設計序，本程式除了退出以外，並沒有執行其他的功能。目的 退出並向linux核心返回乙個狀態碼的簡單程式 輸入 無 輸出 返回乙個狀態碼.在執行程式後可通過輸入echo 來讀取狀態碼 變數 eax儲存系統呼叫號 ebx儲存返回狀態 section data section text g...
微機匯程式設計序
又是自學的一學期，呵呵。學到最後也就知道零星半點指令吧，複雜的程式可能還是不怎麼會寫，熟練當然也不敵c了，但是彙編之於嵌入式，往上走肯定少不了遇到，學好還是必要的！此次僅作入門吧。今日所學，明日之用。1 統計正負零的個數 datas segment array db 1,2,1,0,2,0,2,4,...
匯程式設計序呼叫c程式
首先是匯程式設計序，還是前面的例子，只是加了2行程式 extern main 說明這個函式從外面程式獲得 section data charact db a section text global start start mov ecx,charact push ecx call usestack ...