Source code of the third example that can be found at http://codeofhonour.blogspot.it/2014/10/memories-from-cuda-constant-memory-ii.html

device_transfer_data.cu

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


int host_x[4] = {1, 2, 3, 4};
__device__ int dev_x[4];

__global__ void kernel(int *d_var) { d_var[threadIdx.x] += 10; }

int main(void)
{
	/* Declarations */
	int data_size = 4 * sizeof(int);
	int *address;

	/* Allocate memory for `dev_x` using `cudaMallocHost` */
	checkCudaErrors(cudaMallocHost((void**) &dev_x, data_size, cudaHostAllocWriteCombined));

	/* Get the address of the `__device__` variable `dev_x` */
	checkCudaErrors(cudaGetSymbolAddress((void**)&address, dev_x));	

	/* Transfer data to `address` on the device */
	checkCudaErrors(cudaMemcpy(address, host_x, data_size,cudaMemcpyHostToDevice));
	
	/* Launch the kernel for `address` */
	kernel<<<1,4>>>(address);

	checkCudaErrors(cudaDeviceSynchronize());
	getLastCudaError("wtf!");

	checkCudaErrors(cudaMemcpyFromSymbol(host_x, dev_x, data_size, 0, cudaMemcpyDeviceToHost));

	for (int i=0; i< 4; i++){
		printf("%d\n", host_x[i]);
	}
	return 0;
}

why we need line 18? I think if we don't allocate memory on host for dev_x, it still ok.