How to get SM number?

I have a question. How to get current SM no. in kernel code? For example:

global test (int* A){

A[blockIdx.x]= SM_num; //SM_num is SM number executing current thread.

}

How can I get SM_num? Thanks.

http://forums.nvidia.com/index.php?showtopic=186669