Templated kernel and device function

devicefunction.h

#ifndef DEVICEFUNCTION_H_
#define DEVICEFUNCTION_H_

__device__ void RestoreThreadState(SimState *d_state, GPUThreadStates *tstates,
  	PhotonStructGPU *photon, UINT64 *rnd_x, UINT32 *rnd_a,
		UINT32 *is_active) {
	UINT32 tid = blockIdx.x * NUM_THREADS_PER_BLOCK + threadIdx.x;

	*rnd_x = d_state->x[tid];
	int t = d_state->a[tid];

	*rnd_a = t;

	photon->x = tstates->photon_x[tid];
	photon->y = tstates->photon_y[tid];
	photon->z = tstates->photon_z[tid];
	photon->ux = tstates->photon_ux[tid];
	photon->uy = tstates->photon_uy[tid];
	photon->uz = tstates->photon_uz[tid];
	photon->w = tstates->photon_w[tid];
	photon->sleft = tstates->photon_sleft[tid];
	photon->layer = tstates->photon_layer[tid];

	*is_active = tstates->is_active[tid];
}

#endif /* DEVICEFUNCTION_H_ */

devicekernel.h

#ifndef DEVICEKERNEL_H_
#define DEVICEKERNEL_H_


template<int ignoreAdetection>
__global__ void MCMLKernel(SimState d_state, GPUThreadStates tstates) {
  // photon structure stored in registers
	PhotonStructGPU photon;

	// random number seeds
	UINT64 rnd_x;
	UINT32 rnd_a;

	// Flag to indicate if this thread is active
	UINT32 is_active;

	// Restore the thread state from global memory.
	RestoreThreadState(&d_state, &tstates, &photon, &rnd_x, &rnd_a, &is_active);
}



#endif /* DEVICEKERNEL_H_ */

main.cu

#include "types.h"
#include "devicefunction.h"
#include "devicekernel.h"

int main(int, char**) {
	SimState ss = {};
	GPUThreadStates ts = {};
	MCMLKernel<4> <<<1, 1>>>(ss, ts);
	return 0;
}

types.h

#ifndef TYPES_H_
#define TYPES_H_

typedef unsigned long UINT64;
typedef unsigned int UINT32;

struct SimState {
  int* x;
	int* y;
	int* z;
	int* uy;
	int* uz;
	int* w;
	int* sleft;
	int* layer;
	int* a;
};
struct GPUThreadStates {
	int* photon_x;
	int* photon_y;
	int* photon_z;
	int* photon_ux;
	int* photon_uy;
	int* photon_uz;
	int* photon_w;
	int* photon_sleft;
	int* photon_layer;
	int* is_active;
};

struct PhotonStructGPU {
	int x, y, z, uy, uz, w, sleft, layer, a, ux;
};

#define NUM_THREADS_PER_BLOCK 1

#endif /* TYPES_H_ */