hip_kernel_utils.h File Reference

#include <hip/hip_runtime.h>
#include "gromacs/utility/basedefinitions.h"

Include dependency graph for hip_kernel_utils.h:

Description

HIP device util functions (callable from GPU kernels).

Author

Szilard Pall pall..nosp@m.szil.nosp@m.ard@g.nosp@m.mail.nosp@m..com

Paul Bauer paul..nosp@m.baue.nosp@m.r.q@g.nosp@m.mail.nosp@m..com

Macros
#define	LAUNCH_BOUNDS_EXACT(WORK_GROUP_SIZE, WAVES_PER_EU)
#define	LAUNCH_BOUNDS_EXACT_SINGLE(WORK_GROUP_SIZE)   __attribute__((amdgpu_flat_work_group_size(WORK_GROUP_SIZE, WORK_GROUP_SIZE)))
#define	GMX_HIP_MAX_BLOCKS_PER_MP   16
#define	GMX_HIP_MAX_THREADS_PER_MP   1024

Functions
template<typename T >
__device__ __forceinline__ T	LDG (const T *ptr)
template<typename T >
static __forceinline__ __device__ T	fetchFromParamLookupTable (const T *d_ptr, const hipTextureObject_t texObj, int index)
	Fetch the value by index from the parameter lookup table. More...

Macro Definition Documentation

#define LAUNCH_BOUNDS_EXACT	(		WORK_GROUP_SIZE,
			WAVES_PER_EU
	)

Value:

__attribute__((amdgpu_flat_work_group_size(WORK_GROUP_SIZE, WORK_GROUP_SIZE), \
                   amdgpu_waves_per_eu(WAVES_PER_EU, WAVES_PER_EU)))

Function Documentation

template<typename T >

static __forceinline__ __device__ T fetchFromParamLookupTable ( const T *  d_ptr, const hipTextureObject_t  texObj, int  index  )

static

Fetch the value by index from the parameter lookup table.

Depending on what is supported, it fetches parameters either using direct load or texture objects.

Template Parameters

T	Raw data type

Parameters

[in]	d_ptr	Device pointer to the raw table memory
[in]	texObj	Table texture object
[in]	index	Non-negative element index

Returns

The value from the table at index

template<typename T >

__device__ __forceinline__ T LDG

(

const T *

ptr

)

Load directly or using __ldg() when supported.