Probe the CUDA special registers %smid and %nsmid.

smid.cu

#include <stdio.h>

//
//
//

#define DEVICE_INTRINSIC_QUALIFIERS   __device__ __forceinline__

DEVICE_INTRINSIC_QUALIFIERS
unsigned int
smid()
{
  unsigned int r;

  asm("mov.u32 %0, %%smid;" : "=r"(r));

  return r;
}

DEVICE_INTRINSIC_QUALIFIERS
unsigned int
nsmid()
{
#if (__CUDA_ARCH__ >= 200)
  unsigned int r;

  asm("mov.u32 %0, %%nsmid;" : "=r"(r));

  return r;
#else
  return 30;
#endif
}


//
//
//

__device__ int g_nsmid[1];
__device__ int g_smid[64] = {
  -1, -1, -1, -1, -1, -1, -1, -1,
  -1, -1, -1, -1, -1, -1, -1, -1,
  -1, -1, -1, -1, -1, -1, -1, -1,
  -1, -1, -1, -1, -1, -1, -1, -1,
  -1, -1, -1, -1, -1, -1, -1, -1,
  -1, -1, -1, -1, -1, -1, -1, -1,
  -1, -1, -1, -1, -1, -1, -1, -1,
  -1, -1, -1, -1, -1, -1, -1, -1
};

//
//
//

__global__
void
smidTest()
{
  const int n = nsmid();
  const int s = smid();

  g_nsmid[0] = n;
  g_smid[s]  = s;
}

//
//
//

int main(int argc, char **argv)
{
  cudaError_t err;
  int         device = (argc == 1) ? 0 : atoi(argv[1]);

  cudaDeviceProp props;
  err = cudaGetDeviceProperties(&props,device);

  if (err) 
    return -1;

  // if (props.major < 2) {
  //   printf("%s = sm_%d%d\n",props.name,props.major,props.minor);
  //   return -1;
  // }

  cudaSetDevice(device);
                            
  //
  // LAUNCH KERNEL
  //
  
  smidTest<<<props.multiProcessorCount,1,props.sharedMemPerBlock-384>>>();

  cudaDeviceSynchronize();

  //
  // LOOK AT RESULTS
  //

  int h_nsmid[1];
  int h_smid[48];

  cudaMemcpyFromSymbol(h_nsmid,g_nsmid,sizeof(h_nsmid));
  cudaMemcpyFromSymbol(h_smid, g_smid, sizeof(h_smid));

  //
  //
  //

  printf("%s (%2d) [ %2d",
         props.name,
         (h_nsmid[0] == 30) ? props.multiProcessorCount : h_nsmid[0],
         h_smid[0]);

  int last = 0;
  
  for (int ii=1; ii<h_nsmid[0]; ii++)
    {
      if (h_smid[ii] != -1)
        last = ii;
    }

  for (int ii=1; ii<=last; ii++)
    {
      const int s = h_smid[ii];
      if (s == -1)
        printf(", --");
      else
        printf(", %2d",s);
    }

  printf(" ]\n");

  //
  //
  //

  return 0;
}

Compiled with:

nvcc -m 32 
-gencode arch=compute_11,code=sm_11 
-gencode arch=compute_20,code=sm_20, 
-gencode arch=compute_30,code=sm_30, 
-Xptxas=-v -o smid smid.cu

Some results:

Tesla   K20c    (13) [  0,  1,  2,  3,  4,  5,  6,  7,  8,  9, 10, 11, 12 ]
GeForce GT 240  (12) [  0,  1,  2, --,  4,  5,  6, --,  8,  9, 10, --, 12, 13, 14, ... ]
GeForce GTX 680 ( 8) [  0,  1,  2,  3,  4,  5,  6,  7 ]
GeForce GT 545  ( 3) [  0,  1,  2 ]
GeForce 9400 GT ( 4) [  0,  1, --, --,  4,  5, ... ]