mtgp32-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 "mtgp32-fast.h"

Data Structures

struct mtgp32_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   23209

#define N   726

#define THREAD_NUM   512

#define LARGE_SIZE   (THREAD_NUM * 3)

#define BLOCK_NUM   32

#define TBL_SIZE   16

Functions

__device__ uint32_t para_rec (uint32_t X1, uint32_t X2, uint32_t Y, int bid)

The function of the recursion formula calculation.

__device__ uint32_t temper (uint32_t V, uint32_t T, int bid)

The tempering function.

__device__ uint32_t temper_single (uint32_t V, uint32_t T, int bid)

The tempering and converting function.

__device__ void status_read (uint32_t status[LARGE_SIZE], const mtgp32_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 (mtgp32_kernel_status_t *d_status, const uint32_t status[LARGE_SIZE], int bid, int tid)

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

__global__ void mtgp32_uint32_kernel (mtgp32_kernel_status_t *d_status, uint32_t *d_data, int size)

kernel function.

__global__ void mtgp32_single_kernel (mtgp32_kernel_status_t *d_status, uint32_t *d_data, int size)

kernel function.

void make_constant (const mtgp32_params_fast_t params[])

This function sets constants in device memory.

void make_kernel_data (mtgp32_kernel_status_t *d_status, mtgp32_params_fast_t params[])

This function initializes kernel I/O data.

void print_float_array (const float array[], int size, int block)

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

void print_uint32_array (uint32_t array[], int size, int block)

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

void make_uint32_random (mtgp32_kernel_status_t *d_status, int num_data)

host function.

void make_single_random (mtgp32_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 single_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 mask = 0xff800000

__shared__ uint32_t status [LARGE_SIZE]

Shared memory The generator's internal status vector.

Detailed Description

Sample Program for CUDA 2.2.

MTGP32-23209 This program generates 32-bit unsigned integers. The period of generated integers is 2²³²⁰⁹-1.

This also generates single precision floating point numbers uniformly distributed in the range [1, 2). (float 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

#define LARGE_SIZE (THREAD_NUM * 3)

#define MEXP 23209

#define N 726

#define TBL_SIZE 16

#define THREAD_NUM 512

Function Documentation

int main	(	int	argc,
		char **	argv
	)

References BLOCK_NUM, LARGE_SIZE, make_constant(), make_kernel_data(), make_single_random(), make_uint32_random(), and mtgp32_params_fast_23209.

void make_constant ( const mtgp32_params_fast_t params[] )

This function sets constants in device memory.

Parameters:

[in] params input, MTGP32 parameters.

References BLOCK_NUM, MTGP32_PARAMS_FAST_T::flt_tmp_tbl, mask, MTGP32_PARAMS_FAST_T::mask, param_tbl, MTGP32_PARAMS_FAST_T::pos, pos_tbl, MTGP32_PARAMS_FAST_T::sh1, sh1_tbl, MTGP32_PARAMS_FAST_T::sh2, sh2_tbl, single_temper_tbl, MTGP32_PARAMS_FAST_T::tbl, TBL_SIZE, temper_tbl, and MTGP32_PARAMS_FAST_T::tmp_tbl.

void make_kernel_data	(	mtgp32_kernel_status_t *	d_status,
		mtgp32_params_fast_t	params[]
	)

This function initializes kernel I/O data.

Parameters:

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

References BLOCK_NUM, mtgp32_init_state(), and status.

void make_single_random	(	mtgp32_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_float_array().

Referenced by main().

void make_uint32_random	(	mtgp32_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_uint32_array().

Referenced by main().

__global__ void mtgp32_single_kernel	(	mtgp32_kernel_status_t *	d_status,
		uint32_t *	d_data,
		int	size
	)

kernel function.

This function generates single precision floating point numbers in d_data.

Parameters:

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

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

__global__ void mtgp32_uint32_kernel	(	mtgp32_kernel_status_t *	d_status,
		uint32_t *	d_data,
		int	size
	)

kernel function.

This function generates 32-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 LARGE_SIZE, mask, N, para_rec(), param_tbl, pos_tbl, sh1_tbl, sh2_tbl, status, status_read(), status_write(), temper(), and THREAD_NUM.

__device__ uint32_t para_rec	(	uint32_t	X1,
		uint32_t	X2,
		uint32_t	Y,
		int	bid
	)

The function of the recursion formula calculation.

Parameters:

`[in]`	X1	the farthest part of state array.
`[in]`	X2	the second farthest part of state array.
`[in]`	Y	a part of state array.
`[in]`	bid	block id.

Returns:: output

References mask, param_tbl, sh1_tbl, and sh2_tbl.

void print_float_array	(	const float	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_single_random().

void print_uint32_array	(	uint32_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_uint32_random().

__device__ void status_read	(	uint32_t	status[LARGE_SIZE],
		const mtgp32_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, mtgp32_kernel_status_t::status, status, and THREAD_NUM.

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

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

Parameters:

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

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

__device__ uint32_t temper	(	uint32_t	V,
		uint32_t	T,
		int	bid
	)

The tempering function.

Parameters:

`[in]`	V	the output value should be tempered.
`[in]`	T	the tempering helper value.
`[in]`	bid	block id.

Returns:: the tempered value.

References temper_tbl.

__device__ uint32_t temper_single	(	uint32_t	V,
		uint32_t	T,
		int	bid
	)

The tempering and converting function.

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

Parameters:

`[in]`	V	the output value should be tempered.
`[in]`	T	the tempering helper value.
`[in]`	bid	block id.

Returns:: the tempered and converted value.

References single_temper_tbl.

Referenced by mtgp32_single_kernel().

Variable Documentation

__constant__ uint32_t mask = 0xff800000

Referenced by make_constant(), mtgp32_uint32_kernel(), and para_rec().

__constant__ uint32_t param_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 single_temper_tbl[BLOCK_NUM][TBL_SIZE]

Referenced by make_constant(), and temper_single().

__shared__ uint32_t status[LARGE_SIZE]

Shared memory The generator's internal status vector.

__constant__ uint32_t temper_tbl[BLOCK_NUM][TBL_SIZE]


Data Structures
struct	mtgp32_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 23209
#define	N 726
#define	THREAD_NUM 512
#define	LARGE_SIZE (THREAD_NUM * 3)
#define	BLOCK_NUM 32
#define	TBL_SIZE 16
Functions
__device__ uint32_t	para_rec (uint32_t X1, uint32_t X2, uint32_t Y, int bid)
	The function of the recursion formula calculation.
__device__ uint32_t	temper (uint32_t V, uint32_t T, int bid)
	The tempering function.
__device__ uint32_t	temper_single (uint32_t V, uint32_t T, int bid)
	The tempering and converting function.
__device__ void	status_read (uint32_t status[LARGE_SIZE], const mtgp32_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 (mtgp32_kernel_status_t *d_status, const uint32_t status[LARGE_SIZE], int bid, int tid)
	Read the internal state vector from shared memory, and write them into kernel I/O data.
__global__ void	mtgp32_uint32_kernel (mtgp32_kernel_status_t d_status, uint32_t d_data, int size)
	kernel function.
__global__ void	mtgp32_single_kernel (mtgp32_kernel_status_t d_status, uint32_t d_data, int size)
	kernel function.
void	make_constant (const mtgp32_params_fast_t params[])
	This function sets constants in device memory.
void	make_kernel_data (mtgp32_kernel_status_t *d_status, mtgp32_params_fast_t params[])
	This function initializes kernel I/O data.
void	print_float_array (const float array[], int size, int block)
	This function is used to compare the outputs with C program's.
void	print_uint32_array (uint32_t array[], int size, int block)
	This function is used to compare the outputs with C program's.
void	make_uint32_random (mtgp32_kernel_status_t *d_status, int num_data)
	host function.
void	make_single_random (mtgp32_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	single_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	mask = 0xff800000
__shared__ uint32_t	status [LARGE_SIZE]
	Shared memory The generator's internal status vector.