$ cat t349.cu
#include <stdio.h>
#include <stdlib.h>
#include <pthread.h>

#define NUM_THREADS 100

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)

__device__ int dcount = 0;
__global__ void kernel()
{
    atomicAdd(&dcount, 1);
}

void *thread_func(void *id)
{
    int my_id = *(int *)id;
    kernel<<<1,1>>>();
    cudaDeviceSynchronize();
    cudaCheckErrors("kernel fail");
    printf("*%d", my_id);
    *(int *)id = -1;
    return NULL;
}

int main(void)
{

    // Make thread
    pthread_t pthread[NUM_THREADS];
    int tids[NUM_THREADS];

    // Create and excute pthread
    int tid = 0;
    for(int i=0; i<NUM_THREADS; i++) {
        tids[i] = tid++;
        pthread_create(&pthread[i], NULL, thread_func, (void *)(tids+i));
    }

    // Join pthread
    for(int i=0; i<NUM_THREADS; i++) {
        pthread_join(pthread[i], NULL);
    }

    int hcount = 0;
    cudaMemcpyFromSymbol(&hcount, dcount, sizeof(int));
    for (int i =0; i<NUM_THREADS; i++)
      if (tids[i] != -1) {printf("thread fail at id %d\n", i); return 1;}
    printf("\nresult = %d\n", hcount);

    return 0;
}
$ nvcc -arch=sm_20 -o t349 t349.cu
$ ./t349
*7*2*20*8*5*18*11*1*16*10*22*24*0*6*25*15*23*3*27*19*29*14*30*12*17*9*31*13*4*32*21*33*26*35*28*37*50*54*62*66*71*77*38*45*51*81*67*97*75*68*99*94*52*86*79*53*39*34*55*47*56*58*60*61*40*59*41*65*63*73*64*72*69*70*42*74*36*76*78*43*44*80*46*82*83*84*85*87*88*90*89*93*91*92*95*49*98*57*48*96
result = 100
$