mtgp64-cuda.cu File Reference

Sample Program for CUDA 2.2. More...

#include <stdio.h>
#include <cutil.h>
#include <stdint.h>
#include <inttypes.h>
#include <errno.h>
#include <stdlib.h>
#include "mtgp64-fast.h"

Data Structures
struct	mtgp64_kernel_status_t
	kernel I/O This structure must be initialized before first use. More...
Defines
#define	__STDC_FORMAT_MACROS 1
#define	__STDC_CONSTANT_MACROS 1
#define	MEXP 44497
#define	N 696
#define	THREAD_NUM 512
#define	LARGE_SIZE (THREAD_NUM * 3)
#define	BLOCK_NUM 32
#define	TBL_SIZE 16
Functions
__device__ void	para_rec (uint32_t RH, uint32_t RL, uint32_t X1H, uint32_t X1L, uint32_t X2H, uint32_t X2L, uint32_t YH, uint32_t YL, int bid)
	The function of the recursion formula calculation.
__device__ uint64_t	temper (uint32_t VH, uint32_t VL, uint32_t TL, int bid)
	The tempering function.
__device__ uint64_t	temper_double (uint32_t VH, uint32_t VL, uint32_t TL, int bid)
	The tempering and converting function.
__device__ void	status_read (uint32_t status[2][LARGE_SIZE], const mtgp64_kernel_status_t *d_status, int bid, int tid)
	Read the internal state vector from kernel I/O data, and put them into shared memory.
__device__ void	status_write (mtgp64_kernel_status_t *d_status, const uint32_t status[2][LARGE_SIZE], int bid, int tid)
	Read the internal state vector from shared memory, and write them into kernel I/O data.
__global__ void	mtgp64_uint64_kernel (mtgp64_kernel_status_t d_status, uint64_t d_data, int size)
	kernel function.
__global__ void	mtgp64_double_kernel (mtgp64_kernel_status_t d_status, uint64_t d_data, int size)
	kernel function.
void	make_constant (const mtgp64_params_fast_t params[])
	This function sets constants in device memory.
void	make_kernel_data (mtgp64_kernel_status_t *d_status, mtgp64_params_fast_t params[])
	This function initializes kernel I/O data.
void	print_double_array (const double array[], int size, int block)
	This function is used to compare the outputs with C program's.
void	print_uint64_array (uint64_t array[], int size, int block)
	This function is used to compare the outputs with C program's.
void	make_uint64_random (mtgp64_kernel_status_t *d_status, int num_data)
	host function.
void	make_double_random (mtgp64_kernel_status_t *d_status, int num_data)
	host function.
int	main (int argc, char **argv)
Variables
__constant__ uint32_t	param_tbl [BLOCK_NUM][TBL_SIZE]
__constant__ uint32_t	temper_tbl [BLOCK_NUM][TBL_SIZE]
__constant__ uint32_t	double_temper_tbl [BLOCK_NUM][TBL_SIZE]
__constant__ uint32_t	pos_tbl [BLOCK_NUM]
__constant__ uint32_t	sh1_tbl [BLOCK_NUM]
__constant__ uint32_t	sh2_tbl [BLOCK_NUM]
__constant__ uint32_t	high_mask = 0xffff8000
__constant__ uint32_t	low_mask = 0x00000000
__shared__ uint32_t	status [2][LARGE_SIZE]
	Shared memory The generator's internal status vector.

Detailed Description

Sample Program for CUDA 2.2.

MTGP64-44497 This program generates 64-bit unsigned integers. The period of generated integers is 2⁴⁴⁴⁹⁷-1.

This also generates double precision floating point numbers uniformly distributed in the range [1, 2). (double r; 1.0 <= r < 2.0)

Author:: Mutsuo Saito (Hiroshima University); Makoto Matsumoto (Hiroshima University)

The new BSD License is applied to this software, see LICENSE.txt

Define Documentation

#define __STDC_CONSTANT_MACROS 1

#define __STDC_FORMAT_MACROS 1

#define BLOCK_NUM 32

Referenced by main(), make_constant(), make_double_random(), make_kernel_data(), make_single_random(), make_uint32_random(), and make_uint64_random().

#define LARGE_SIZE (THREAD_NUM * 3)

Referenced by main(), mtgp32_single_kernel(), mtgp32_uint32_kernel(), mtgp64_double_kernel(), and mtgp64_uint64_kernel().

#define MEXP 44497

#define N 696

Referenced by mtgp32_single_kernel(), mtgp32_uint32_kernel(), mtgp64_double_kernel(), mtgp64_uint64_kernel(), status_read(), and status_write().

#define TBL_SIZE 16

Referenced by make_constant().

#define THREAD_NUM 512

Referenced by mtgp32_single_kernel(), mtgp32_uint32_kernel(), mtgp64_double_kernel(), mtgp64_uint64_kernel(), status_read(), and status_write().

Function Documentation

int main	(	int	argc,
		char **	argv
	)

References BLOCK_NUM, LARGE_SIZE, make_constant(), make_double_random(), make_kernel_data(), make_uint64_random(), and mtgp64_params_fast_44497.

void make_constant ( const mtgp64_params_fast_t params[] )

This function sets constants in device memory.

Parameters:

[in] params input, MTGP64 parameters.

References BLOCK_NUM, MTGP64_PARAMS_FAST_T::dbl_tmp_tbl, double_temper_tbl, high_mask, low_mask, MTGP64_PARAMS_FAST_T::mask, param_tbl, MTGP64_PARAMS_FAST_T::pos, pos_tbl, MTGP64_PARAMS_FAST_T::sh1, sh1_tbl, MTGP64_PARAMS_FAST_T::sh2, sh2_tbl, MTGP64_PARAMS_FAST_T::tbl, TBL_SIZE, temper_tbl, and MTGP64_PARAMS_FAST_T::tmp_tbl.

Referenced by main().

void make_double_random	(	mtgp64_kernel_status_t *	d_status,
		int	num_data
	)

host function.

This function calls corresponding kernel function.

Parameters:

`[in]`	d_status	kernel I/O data.
`[in]`	num_data	number of data to be generated.

References BLOCK_NUM, and print_double_array().

Referenced by main().

void make_kernel_data	(	mtgp64_kernel_status_t *	d_status,
		mtgp64_params_fast_t	params[]
	)

This function initializes kernel I/O data.

Parameters:

`[out]`	d_status	output kernel I/O data.
`[in]`	params	MTGP64 parameters. needed for the initialization.

References BLOCK_NUM, mtgp64_init_state(), and status.

Referenced by main().

void make_uint64_random	(	mtgp64_kernel_status_t *	d_status,
		int	num_data
	)

host function.

This function calls corresponding kernel function.

Parameters:

`[in]`	d_status	kernel I/O data.
`[in]`	num_data	number of data to be generated.

References BLOCK_NUM, and print_uint64_array().

Referenced by main().

__global__ void mtgp64_double_kernel	(	mtgp64_kernel_status_t *	d_status,
		uint64_t *	d_data,
		int	size
	)

kernel function.

This function generates double precision floating point numbers in d_data.

Parameters:

`[in,out]`	d_status	kernel I/O data
`[out]`	d_data	output. IEEE double precision format.
`[in]`	size	number of output data requested.

References LARGE_SIZE, N, para_rec(), pos_tbl, status, status_read(), status_write(), temper_double(), and THREAD_NUM.

__global__ void mtgp64_uint64_kernel	(	mtgp64_kernel_status_t *	d_status,
		uint64_t *	d_data,
		int	size
	)

kernel function.

This function generates 64-bit unsigned integers in d_data

Parameters:

`[in,out]`	d_status	kernel I/O data
`[out]`	d_data	output
`[in]`	size	number of output data requested.

References high_mask, LARGE_SIZE, low_mask, N, para_rec(), param_tbl, pos_tbl, sh1_tbl, sh2_tbl, status, status_read(), status_write(), temper(), and THREAD_NUM.

__device__ void para_rec	(	uint32_t *	RH,
		uint32_t *	RL,
		uint32_t	X1H,
		uint32_t	X1L,
		uint32_t	X2H,
		uint32_t	X2L,
		uint32_t	YH,
		uint32_t	YL,
		int	bid
	)

The function of the recursion formula calculation.

Parameters:

`[out]`	RH	32-bit MSBs of output
`[out]`	RL	32-bit LSBs of output
`[in]`	X1H	MSBs of the farthest part of state array.
`[in]`	X1L	LSBs of the farthest part of state array.
`[in]`	X2H	MSBs of the second farthest part of state array.
`[in]`	X2L	LSBs of the second farthest part of state array.
`[in]`	YH	MSBs of a part of state array.
`[in]`	YL	LSBs of a part of state array.
`[in]`	bid	block id.

References high_mask, low_mask, param_tbl, sh1_tbl, and sh2_tbl.

Referenced by mtgp32_single_kernel(), mtgp32_uint32_kernel(), mtgp64_double_kernel(), and mtgp64_uint64_kernel().

void print_double_array	(	const double	array[],
		int	size,
		int	block
	)

This function is used to compare the outputs with C program's.

Parameters:

`[in]`	array	data to be printed.
`[in]`	size	size of array.
`[in]`	block	number of blocks.

Referenced by make_double_random().

void print_uint64_array	(	uint64_t	array[],
		int	size,
		int	block
	)

This function is used to compare the outputs with C program's.

Parameters:

`[in]`	array	data to be printed.
`[in]`	size	size of array.
`[in]`	block	number of blocks.

Referenced by make_uint64_random().

__device__ void status_read	(	uint32_t	status[2][LARGE_SIZE],
		const mtgp64_kernel_status_t *	d_status,
		int	bid,
		int	tid
	)

Read the internal state vector from kernel I/O data, and put them into shared memory.

Parameters:

`[out]`	status	shared memory.
`[in]`	d_status	kernel I/O data
`[in]`	bid	block id
`[in]`	tid	thread id

References N, status, mtgp64_kernel_status_t::status, and THREAD_NUM.

Referenced by mtgp32_single_kernel(), mtgp32_uint32_kernel(), mtgp64_double_kernel(), and mtgp64_uint64_kernel().

__device__ void status_write	(	mtgp64_kernel_status_t *	d_status,
		const uint32_t	status[2][LARGE_SIZE],
		int	bid,
		int	tid
	)

Read the internal state vector from shared memory, and write them into kernel I/O data.

Parameters:

`[out]`	status	shared memory.
`[in]`	d_status	kernel I/O data
`[in]`	bid	block id
`[in]`	tid	thread id

References N, mtgp64_kernel_status_t::status, status, and THREAD_NUM.

Referenced by mtgp32_single_kernel(), mtgp32_uint32_kernel(), mtgp64_double_kernel(), and mtgp64_uint64_kernel().

__device__ uint64_t temper	(	uint32_t	VH,
		uint32_t	VL,
		uint32_t	TL,
		int	bid
	)

The tempering function.

Parameters:

`[in]`	VH	MSBs of the output value should be tempered.
`[in]`	VL	LSBs of the output value should be tempered.
`[in]`	TL	LSBs of the tempering helper value.
`[in]`	bid	block id.

Returns:: [in] the tempered value.

References temper_tbl.

Referenced by mtgp32_uint32_kernel(), and mtgp64_uint64_kernel().

__device__ uint64_t temper_double	(	uint32_t	VH,
		uint32_t	VL,
		uint32_t	TL,
		int	bid
	)

The tempering and converting function.

By using the preset-ted table, converting to IEEE format and tempering are done simultaneously.

Parameters:

`[in]`	VH	MSBs of the output value should be tempered.
`[in]`	VL	LSBs of the output value should be tempered.
`[in]`	TL	LSBs of the tempering helper value.
`[in]`	bid	block id.