Source code of the last example at http://codeofhonour.blogspot.it/2014/10/static-allocation-of-device-vars.html

device_static_alloc_pinned.cu

#include <stdio.h>

#include <cuda_runtime.h>
#include "helper_cuda.h"

#define NS 3

/**
 * Statically allocated memory on the device
 */__device__ float dev_static[NS] = { 10.0, -20.0, 235.0 };

/**
 * Adds to a given vector `dx`,
 * the vector `dev_static`
 */
 __global__ void kernel(float *dx) {
	int tid = threadIdx.x + blockIdx.x * blockDim.x;
	if (tid < NS)
		dx[tid] += dev_static[tid];
}

int main(void) {

	cudaDeviceProp prop;
	int whichDevice;
	float * hst_p = NULL;
	float * dev_address = NULL;
	int size = NS * sizeof(float);
	hst_p = (float *) malloc(size);

	checkCudaErrors(cudaGetDevice(&whichDevice));
	checkCudaErrors(cudaGetDeviceProperties(&prop, whichDevice));

	if (prop.canMapHostMemory != 1) {
		fprintf(stderr, "Device cannot map memory!\n");
		return 1;
	}

	checkCudaErrors(cudaSetDeviceFlags(cudaDeviceMapHost));

	for (int i = 0; i < NS; i++) {
		hst_p[i] = (float) (i + 1);
	}

	checkCudaErrors(cudaHostAlloc(&hst_p, size, cudaHostAllocMapped));

	checkCudaErrors(cudaHostGetDevicePointer(&dev_address, hst_p, 0));

	kernel<<<1, NS>>>(dev_address);
	/*
	 * The following line is necessary for the host
	 * to be able to "see" the changes that have been done
	 * on `host_p`
	 */
	checkCudaErrors(cudaDeviceSynchronize());

	for (int i = 0; i < NS; i++) {
		printf("host_p[%d] = %g\n", i, hst_p[i]);
	}

	if (hst_p != NULL)
		checkCudaErrors(cudaFreeHost(hst_p));

	return 0;
}