Ben Sander
@@ -0,0 +1 @@
|
||||
|
||||
@@ -0,0 +1 @@
|
||||
|
||||
@@ -0,0 +1,556 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
/**
|
||||
* @file hcc_detail/hip_runtime.h
|
||||
*
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
|
||||
//---
|
||||
// Top part of file can be compiled with any compiler
|
||||
|
||||
|
||||
#include <cstring>
|
||||
#include <cmath>
|
||||
#include <string.h>
|
||||
#include <stddef.h>
|
||||
|
||||
|
||||
#define CUDA_SUCCESS hipSuccess
|
||||
|
||||
#include <hip_runtime_api.h>
|
||||
|
||||
//---
|
||||
// Remainder of this file only compiles with HCC
|
||||
#ifdef __HCC__
|
||||
#include <hc.hpp>
|
||||
#include <grid_launch.h>
|
||||
|
||||
//TODO-HCC-GL - change this to typedef.
|
||||
//typedef grid_launch_parm hipLaunchParm ;
|
||||
#define hipLaunchParm grid_launch_parm
|
||||
|
||||
#include <hcc_detail/hip_texture.h>
|
||||
#include <hcc_detail/host_defines.h>
|
||||
|
||||
// TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define.
|
||||
#if defined (__KALMAR_ACCELERATOR__) && not defined (__HCC_ACCELERATOR__)
|
||||
#define __HCC_ACCELERATOR__ __KALMAR_ACCELERATOR__
|
||||
#endif
|
||||
|
||||
// Feature tests:
|
||||
#if defined(__HCC_ACCELERATOR__) and (__HCC_ACCELERATOR__ != 0)
|
||||
// Device compile and not host compile:
|
||||
|
||||
|
||||
#define __HIP_DEVICE_COMPILE__ 1
|
||||
|
||||
//TODO-HCC enable __HIP_ARCH_HAS_ATOMICS__ when HCC supports these.
|
||||
// 32-bit Atomics:
|
||||
#define __HIP_ARCH_HAS_GLOBAL_INT32_ATOMICS__ (1)
|
||||
#define __HIP_ARCH_HAS_GLOBAL_FLOAT_ATOMIC_EXCH__ (1)
|
||||
#define __HIP_ARCH_HAS_SHARED_INT32_ATOMICS__ (0)
|
||||
#define __HIP_ARCH_HAS_SHARED_FLOAT_ATOMIC_EXCH__ (0)
|
||||
#define __HIP_ARCH_HAS_FLOAT_ATOMIC_ADD__ (0)
|
||||
|
||||
// 64-bit Atomics:
|
||||
#define __HIP_ARCH_HAS_GLOBAL_INT64_ATOMICS__ (1)
|
||||
#define __HIP_ARCH_HAS_SHARED_INT64_ATOMICS__ (0)
|
||||
|
||||
// Doubles
|
||||
#define __HIP_ARCH_HAS_DOUBLES__ (1)
|
||||
|
||||
//warp cross-lane operations:
|
||||
#define __HIP_ARCH_HAS_WARP_VOTE__ (1)
|
||||
#define __HIP_ARCH_HAS_WARP_BALLOT__ (1)
|
||||
#define __HIP_ARCH_HAS_WARP_SHUFFLE__ (1)
|
||||
#define __HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ (0)
|
||||
|
||||
//sync
|
||||
#define __HIP_ARCH_HAS_THREAD_FENCE_SYSTEM__ (0)
|
||||
#define __HIP_ARCH_HAS_SYNC_THREAD_EXT__ (0)
|
||||
|
||||
// misc
|
||||
#define __HIP_ARCH_HAS_SURFACE_FUNCS__ (0)
|
||||
#define __HIP_ARCH_HAS_3DGRID__ (1)
|
||||
#define __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ (0)
|
||||
|
||||
#else
|
||||
// Host compile and not device compile:
|
||||
#define __HIP_DEVICE_COMPILE__ 0
|
||||
|
||||
#endif
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
//TODO-HCC this is currently ignored by HCC target of HIP
|
||||
#define __launch_bounds__(requiredMaxThreadsPerBlock, minBlocksPerMultiprocessor)
|
||||
|
||||
// Detect if we are compiling C++ mode or C mode
|
||||
#if defined(__cplusplus)
|
||||
#define __HCC_CPP__
|
||||
#elif defined(__STDC_VERSION__)
|
||||
#define __HCC_C__
|
||||
#endif
|
||||
|
||||
#define clock_t long long int
|
||||
__device__ inline long long int clock64() { return (long long int)hc::__clock_u64(); };
|
||||
__device__ inline clock_t clock() { return (clock_t)hc::__clock_u64(); };
|
||||
|
||||
//atomicAdd()
|
||||
__device__ inline int atomicAdd(int* address, int val)
|
||||
{
|
||||
return hc::atomic_fetch_add(address,val);
|
||||
}
|
||||
__device__ inline unsigned int atomicAdd(unsigned int* address,
|
||||
unsigned int val)
|
||||
{
|
||||
return hc::atomic_fetch_add(address,val);
|
||||
}
|
||||
__device__ inline unsigned long long int atomicAdd(unsigned long long int* address,
|
||||
unsigned long long int val)
|
||||
{
|
||||
return (long long int)hc::atomic_fetch_add((uint64_t*)address,(uint64_t)val);
|
||||
}
|
||||
__device__ inline float atomicAdd(float* address, float val)
|
||||
{
|
||||
return hc::atomic_fetch_add(address,val);
|
||||
}
|
||||
|
||||
//atomicSub()
|
||||
__device__ inline int atomicSub(int* address, int val)
|
||||
{
|
||||
return hc::atomic_fetch_sub(address,val);
|
||||
}
|
||||
__device__ inline unsigned int atomicSub(unsigned int* address,
|
||||
unsigned int val)
|
||||
{
|
||||
return hc::atomic_fetch_sub(address,val);
|
||||
}
|
||||
|
||||
//atomicExch()
|
||||
__device__ inline int atomicExch(int* address, int val)
|
||||
{
|
||||
return hc::atomic_exchange(address,val);
|
||||
}
|
||||
__device__ inline unsigned int atomicExch(unsigned int* address,
|
||||
unsigned int val)
|
||||
{
|
||||
return hc::atomic_exchange(address,val);
|
||||
}
|
||||
__device__ inline unsigned long long int atomicExch(unsigned long long int* address,
|
||||
unsigned long long int val)
|
||||
{
|
||||
return (long long int)hc::atomic_exchange((uint64_t*)address,(uint64_t)val);
|
||||
}
|
||||
__device__ inline float atomicExch(float* address, float val)
|
||||
{
|
||||
return hc::atomic_exchange(address,val);
|
||||
}
|
||||
|
||||
//atomicMin()
|
||||
__device__ inline int atomicMin(int* address, int val)
|
||||
{
|
||||
return hc::atomic_fetch_min(address,val);
|
||||
}
|
||||
__device__ inline unsigned int atomicMin(unsigned int* address,
|
||||
unsigned int val)
|
||||
{
|
||||
return hc::atomic_fetch_min(address,val);
|
||||
}
|
||||
__device__ inline unsigned long long int atomicMin(unsigned long long int* address,
|
||||
unsigned long long int val)
|
||||
{
|
||||
return (long long int)hc::atomic_fetch_min((uint64_t*)address,(uint64_t)val);
|
||||
}
|
||||
|
||||
//atomicMax()
|
||||
__device__ inline int atomicMax(int* address, int val)
|
||||
{
|
||||
return hc::atomic_fetch_max(address,val);
|
||||
}
|
||||
__device__ inline unsigned int atomicMax(unsigned int* address,
|
||||
unsigned int val)
|
||||
{
|
||||
return hc::atomic_fetch_max(address,val);
|
||||
}
|
||||
__device__ inline unsigned long long int atomicMax(unsigned long long int* address,
|
||||
unsigned long long int val)
|
||||
{
|
||||
return (long long int)hc::atomic_fetch_max((uint64_t*)address,(uint64_t)val);
|
||||
}
|
||||
|
||||
//atomicInc()
|
||||
__device__ inline unsigned int atomicInc(unsigned int* address)
|
||||
{
|
||||
return hc::atomic_fetch_inc(address);
|
||||
}
|
||||
|
||||
//atomicDec()
|
||||
__device__ inline unsigned int atomicDec(unsigned int* address)
|
||||
{
|
||||
return hc::atomic_fetch_dec(address);
|
||||
}
|
||||
|
||||
//atomicCAS()
|
||||
__device__ inline int atomicCAS(int* address, int compare, int val)
|
||||
{
|
||||
hc::atomic_compare_exchange(address,&compare,val);
|
||||
return *address;
|
||||
}
|
||||
__device__ inline unsigned int atomicCAS(unsigned int* address,
|
||||
unsigned int compare,
|
||||
unsigned int val)
|
||||
{
|
||||
hc::atomic_compare_exchange(address,&compare,val);
|
||||
return *address;
|
||||
}
|
||||
__device__ inline unsigned long long int atomicCAS(unsigned long long int* address,
|
||||
unsigned long long int compare,
|
||||
unsigned long long int val)
|
||||
{
|
||||
hc::atomic_compare_exchange((uint64_t*)address,(uint64_t*)&compare,(uint64_t)val);
|
||||
return *address;
|
||||
}
|
||||
|
||||
//atomicAnd()
|
||||
__device__ inline int atomicAnd(int* address, int val)
|
||||
{
|
||||
return hc::atomic_fetch_and(address,val);
|
||||
}
|
||||
__device__ inline unsigned int atomicAnd(unsigned int* address,
|
||||
unsigned int val)
|
||||
{
|
||||
return hc::atomic_fetch_and(address,val);
|
||||
}
|
||||
__device__ inline unsigned long long int atomicAnd(unsigned long long int* address,
|
||||
unsigned long long int val)
|
||||
{
|
||||
return (long long int)hc::atomic_fetch_and((uint64_t*)address,(uint64_t)val);
|
||||
}
|
||||
|
||||
//atomicOr()
|
||||
__device__ inline int atomicOr(int* address, int val)
|
||||
{
|
||||
return hc::atomic_fetch_or(address,val);
|
||||
}
|
||||
__device__ inline unsigned int atomicOr(unsigned int* address,
|
||||
unsigned int val)
|
||||
{
|
||||
return hc::atomic_fetch_or(address,val);
|
||||
}
|
||||
__device__ inline unsigned long long int atomicOr(unsigned long long int* address,
|
||||
unsigned long long int val)
|
||||
{
|
||||
return (long long int)hc::atomic_fetch_or((uint64_t*)address,(uint64_t)val);
|
||||
}
|
||||
|
||||
//atomicXor()
|
||||
__device__ inline int atomicXor(int* address, int val)
|
||||
{
|
||||
return hc::atomic_fetch_xor(address,val);
|
||||
}
|
||||
__device__ inline unsigned int atomicXor(unsigned int* address,
|
||||
unsigned int val)
|
||||
{
|
||||
return hc::atomic_fetch_xor(address,val);
|
||||
}
|
||||
__device__ inline unsigned long long int atomicXor(unsigned long long int* address,
|
||||
unsigned long long int val)
|
||||
{
|
||||
return (long long int)hc::atomic_fetch_xor((uint64_t*)address,(uint64_t)val);
|
||||
}
|
||||
|
||||
#ifdef __HCC__
|
||||
#include <hc.hpp>
|
||||
// integer intrinsic function __poc __clz __ffs __brev
|
||||
__device__ inline unsigned int __popc( unsigned int input)
|
||||
{
|
||||
return hc::__popcount_u32_b32( input);
|
||||
}
|
||||
|
||||
__device__ inline unsigned int __popcll( unsigned long long int input)
|
||||
{
|
||||
return hc::__popcount_u32_b64(input);
|
||||
}
|
||||
|
||||
__device__ inline unsigned int __clz(unsigned int input)
|
||||
{
|
||||
return hc::__firstbit_u32_u32( input);
|
||||
}
|
||||
|
||||
__device__ inline unsigned int __clzll(unsigned long long int input)
|
||||
{
|
||||
return hc::__firstbit_u32_u64( input);
|
||||
}
|
||||
|
||||
__device__ inline unsigned int __clz(int input)
|
||||
{
|
||||
return hc::__firstbit_u32_s32( input);
|
||||
}
|
||||
|
||||
__device__ inline unsigned int __clzll(long long int input)
|
||||
{
|
||||
return hc::__firstbit_u32_s64( input);
|
||||
}
|
||||
|
||||
__device__ inline unsigned int __ffs(unsigned int input)
|
||||
{
|
||||
return hc::__lastbit_u32_u32( input)+1;
|
||||
}
|
||||
|
||||
__device__ inline unsigned int __ffsll(unsigned long long int input)
|
||||
{
|
||||
return hc::__lastbit_u32_u64( input)+1;
|
||||
}
|
||||
|
||||
__device__ inline unsigned int __brev( unsigned int input)
|
||||
{
|
||||
return hc::__bitrev_b32( input);
|
||||
}
|
||||
|
||||
__device__ inline unsigned long long int __brevll( unsigned long long int input)
|
||||
{
|
||||
return hc::__bitrev_b64( input);
|
||||
}
|
||||
|
||||
// warp vote function __all __any __ballot
|
||||
|
||||
__device__ inline int __all( int input)
|
||||
{
|
||||
return hc::__all( input);
|
||||
}
|
||||
|
||||
|
||||
__device__ inline int __any( int input)
|
||||
{
|
||||
return hc::__any( input);
|
||||
}
|
||||
|
||||
|
||||
__device__ inline unsigned long long int __ballot( int input)
|
||||
{
|
||||
return hc::__ballot( input);
|
||||
}
|
||||
|
||||
#endif
|
||||
|
||||
|
||||
|
||||
#ifdef __HCC_ACCELERATOR__
|
||||
#include <hc_math.hpp>
|
||||
// TODO: Choose whether default is precise math or fast math based on compilation flag.
|
||||
using namespace hc::precise_math;
|
||||
|
||||
//TODO: Undo this once min/max functions are supported by hc
|
||||
inline int min(int arg1, int arg2) __attribute((hc,cpu)) { \
|
||||
return (int)(hc::precise_math::fmin((float)arg1, (float)arg2));}
|
||||
inline int max(int arg1, int arg2) __attribute((hc,cpu)) { \
|
||||
return (int)(hc::precise_math::fmax((float)arg1, (float)arg2));}
|
||||
|
||||
|
||||
//TODO - add a couple fast math operations here, the set here will grow :
|
||||
__device__ inline float __log2f(float x) {return hc::fast_math::log2(x); };
|
||||
__device__ inline float __powf(float base, float exponent) {return hc::fast_math::powf(base, exponent); };
|
||||
|
||||
#endif
|
||||
|
||||
|
||||
|
||||
/**
|
||||
* Kernel launching
|
||||
*/
|
||||
#define hipThreadIdx_x (amp_get_local_id(2))
|
||||
#define hipThreadIdx_y (amp_get_local_id(1))
|
||||
#define hipThreadIdx_z (amp_get_local_id(0))
|
||||
|
||||
#define hipBlockIdx_x (hc_get_group_id(2))
|
||||
#define hipBlockIdx_y (hc_get_group_id(1))
|
||||
#define hipBlockIdx_z (hc_get_group_id(0))
|
||||
|
||||
#define hipBlockDim_x (amp_get_local_size(2))
|
||||
#define hipBlockDim_y (amp_get_local_size(1))
|
||||
#define hipBlockDim_z (amp_get_local_size(0))
|
||||
|
||||
#define hipGridDim_x (hc_get_num_groups(2))
|
||||
#define hipGridDim_y (hc_get_num_groups(1))
|
||||
#define hipGridDim_z (hc_get_num_groups(0))
|
||||
|
||||
|
||||
extern int warpSize ;
|
||||
|
||||
|
||||
#define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE)
|
||||
|
||||
|
||||
#if 0
|
||||
#define KALMAR_PFE_BEGIN() \
|
||||
hc::extent<3> ext(lp.gridDim.x, lp.gridDim.y, lp.gridDim.z);\
|
||||
auto __hipExtTile = ext.tile(lp.groupDim.x, lp.groupDim.y, lp.groupDim.z);\
|
||||
__hipExtTile.set_dynamic_group_segment_size(lp.groupMemBytes);\
|
||||
\
|
||||
hc::completion_future cf = hc::parallel_for_each (\
|
||||
*lp.av,\
|
||||
__hipExtTile,\
|
||||
[=] (hc::tiled_index<3> __hipIdx) mutable [[hc]]
|
||||
|
||||
|
||||
|
||||
#define KALMAR_PFE_END \
|
||||
); \
|
||||
if (HIP_LAUNCH_BLOCKING) {\
|
||||
if (HIP_TRACE_API) {\
|
||||
fprintf(stderr, "hiptrace1: HIP_LAUNCH_BLOCKING ...\n");\
|
||||
}\
|
||||
cf.wait(); \
|
||||
if (HIP_TRACE_API) {\
|
||||
fprintf(stderr, "hiptrace1: ...completed.\n");\
|
||||
}\
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
|
||||
#define HIP_KERNEL_NAME(...) __VA_ARGS__
|
||||
|
||||
|
||||
#ifdef __HCC_CPP__
|
||||
hc::accelerator_view *ihipLaunchKernel(hipStream_t stream);
|
||||
|
||||
#if not defined(DISABLE_GRID_LAUNCH)
|
||||
#define hipLaunchKernel(_kernelName, _numBlocks3D, _blockDim3D, _groupMemBytes, _stream, ...) \
|
||||
do {\
|
||||
grid_launch_parm lp;\
|
||||
lp.gridDim.x = _numBlocks3D.x; \
|
||||
lp.gridDim.y = _numBlocks3D.y; \
|
||||
lp.gridDim.z = _numBlocks3D.z; \
|
||||
lp.groupDim.x = _blockDim3D.x; \
|
||||
lp.groupDim.y = _blockDim3D.y; \
|
||||
lp.groupDim.z = _blockDim3D.z; \
|
||||
lp.groupMemBytes = _groupMemBytes;\
|
||||
hc::completion_future cf;\
|
||||
lp.cf = &cf; \
|
||||
lp.av = (ihipLaunchKernel(_stream)); \
|
||||
if (HIP_TRACE_API) {\
|
||||
fprintf(stderr, "hiptrace1: launch '%s' gridDim:[%d.%d.%d] groupDim:[%d.%d.%d] groupMem:+%d stream=%p\n", \
|
||||
#_kernelName, lp.gridDim.z, lp.gridDim.y, lp.gridDim.x, lp.groupDim.z, lp.groupDim.y, lp.groupDim.x, lp.groupMemBytes, (void*)(_stream));\
|
||||
}\
|
||||
_kernelName (lp, __VA_ARGS__);\
|
||||
} while(0)
|
||||
|
||||
#else
|
||||
#warning(DISABLE_GRID_LAUNCH set)
|
||||
|
||||
#define hipLaunchKernel(_kernelName, _numBlocks3D, _blockDim3D, _groupMemBytes, _stream, ...) \
|
||||
do {\
|
||||
grid_launch_parm lp;\
|
||||
lp.gridDim.x = _numBlocks3D.x * _blockDim3D.x;/*Convert from #blocks to #threads*/ \
|
||||
lp.gridDim.y = _numBlocks3D.y * _blockDim3D.y;/*Convert from #blocks to #threads*/ \
|
||||
lp.gridDim.z = _numBlocks3D.z * _blockDim3D.z;/*Convert from #blocks to #threads*/ \
|
||||
lp.groupDim.x = _blockDim3D.x; \
|
||||
lp.groupDim.y = _blockDim3D.y; \
|
||||
lp.groupDim.z = _blockDim3D.z; \
|
||||
lp.groupMemBytes = _groupMemBytes;\
|
||||
hc::completion_future cf;\
|
||||
lp.cf = &cf; \
|
||||
lp.av = (ihipLaunchKernel(_stream)); \
|
||||
if (HIP_TRACE_API) {\
|
||||
fprintf(stderr, "hiptrace1: launch '%s' gridDim:[%d.%d.%d] groupDim:[%d.%d.%d] groupMem:+%d stream=%p\n", \
|
||||
#_kernelName, lp.gridDim.z, lp.gridDim.y, lp.gridDim.x, lp.groupDim.z, lp.groupDim.y, lp.groupDim.x, lp.groupMemBytes, (void*)(_stream));\
|
||||
}\
|
||||
_kernelName (lp, __VA_ARGS__);\
|
||||
} while(0)
|
||||
/*end hipLaunchKernel */
|
||||
#endif
|
||||
|
||||
#elif defined (__HCC_C__)
|
||||
|
||||
//TODO - develop C interface.
|
||||
|
||||
#endif
|
||||
|
||||
|
||||
#if not defined(DISABLE_GRID_LAUNCH)
|
||||
// TODO -In GL these are no-ops and can be removed:
|
||||
// Keep them around for a little while as a fallback.
|
||||
#define KERNELBEGIN
|
||||
#define KERNELEND
|
||||
|
||||
#else
|
||||
|
||||
// TODO-GL:
|
||||
// These wrap the kernel in a PFE loop with macros.
|
||||
// Not required with GL but exist here as a fallback.
|
||||
#define KERNELBEGIN \
|
||||
hc::extent<3> ext(lp.gridDim.x, lp.gridDim.y, lp.gridDim.z);\
|
||||
auto __hipExtTile = ext.tile(lp.groupDim.x, lp.groupDim.y, lp.groupDim.z);\
|
||||
__hipExtTile.set_dynamic_group_segment_size(lp.groupMemBytes);\
|
||||
\
|
||||
hc::completion_future cf = \
|
||||
hc::parallel_for_each (\
|
||||
*lp.av,\
|
||||
__hipExtTile,\
|
||||
[=] (hc::tiled_index<3> __hipIdx) mutable [[hc]] \
|
||||
{
|
||||
|
||||
|
||||
#define KERNELEND \
|
||||
}); \
|
||||
if (HIP_LAUNCH_BLOCKING) {\
|
||||
if (HIP_TRACE_API) {\
|
||||
fprintf(stderr, "hiptrace1: HIP_LAUNCH_BLOCKING ...\n");\
|
||||
}\
|
||||
cf.wait(); \
|
||||
if (HIP_TRACE_API) {\
|
||||
fprintf(stderr, "hiptrace1: ...completed.\n");\
|
||||
}\
|
||||
}
|
||||
|
||||
#endif /*DISABLE_GRID_LAUNCH*/
|
||||
|
||||
|
||||
#endif // __HCC__
|
||||
|
||||
|
||||
/**
|
||||
* @defgroup HIP-ENV HIP Environment Variables
|
||||
* @{
|
||||
*/
|
||||
extern int HIP_PRINT_ENV ; ///< Print all HIP-related environment variables.
|
||||
extern int HIP_TRACE_API; ///< Trace HIP APIs.
|
||||
extern int HIP_LAUNCH_BLOCKING ; ///< Make all HIP APIs host-synchronous
|
||||
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
|
||||
// End doxygen API:
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
|
||||
|
||||
@@ -0,0 +1,871 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
#pragma once
|
||||
|
||||
|
||||
#include <stdint.h>
|
||||
#include <stddef.h>
|
||||
|
||||
#include <hcc_detail/host_defines.h>
|
||||
|
||||
#if defined (__HCC__) && (__hcc_workweek__ < 1602)
|
||||
#error("This version of HIP requires a newer version of HCC.");
|
||||
#endif
|
||||
|
||||
|
||||
// hip_api_hcc.h
|
||||
// Contains C function APIs for HIP runtime.
|
||||
// This file does not use any HCC builtins or special language extensions (-hc mode) ; those functions in hip_hcc.h.
|
||||
|
||||
|
||||
// Structure definitions:
|
||||
#ifdef __cplusplus
|
||||
extern "C" {
|
||||
#endif
|
||||
|
||||
/**
|
||||
* @addtogroup GlobalDefs More
|
||||
* @{
|
||||
*/
|
||||
//! Flags that can be used with hipStreamCreateWithFlags
|
||||
#define hipStreamDefault 0x00 ///< Default stream creation flags. These are used with hipStreamCreate().
|
||||
#define hipStreamNonBlocking 0x01 ///< Stream does not implicitly synchronize with null stream
|
||||
|
||||
|
||||
//! Flags that can be used with hipEventCreateWithFlags:
|
||||
#define hipEventDefault 0x0 ///< Default flags
|
||||
#define hipEventBlockingSync 0x1 ///< Waiting will yield CPU. Power-friendly and usage-friendly but may increase latency.
|
||||
#define hipEventDisableTiming 0x2 ///< Disable event's capability to record timing information. May improve performance.
|
||||
#define hipEventInterprocess 0x4 ///< Event can support IPC. @warning - not supported in HIP.
|
||||
|
||||
|
||||
/**
|
||||
* @warning On AMD devices and recent NVIDIA devices, these hints and controls are ignored.
|
||||
*/
|
||||
typedef enum hipFuncCache {
|
||||
hipFuncCachePreferNone, ///< no preference for shared memory or L1 (default)
|
||||
hipFuncCachePreferShared, ///< prefer larger shared memory and smaller L1 cache
|
||||
hipFuncCachePreferL1, ///< prefer larger L1 cache and smaller shared memory
|
||||
hipFuncCachePreferEqual, ///< prefer equal size L1 cache and shared memory
|
||||
} hipFuncCache;
|
||||
|
||||
|
||||
/**
|
||||
* @warning On AMD devices and recent NVIDIA devices, these hints and controls are ignored.
|
||||
*/
|
||||
typedef enum hipSharedMemConfig {
|
||||
hipSharedMemBankSizeDefault, ///< The compiler selects a device-specific value for the banking.
|
||||
hipSharedMemBankSizeFourByte, ///< Shared mem is banked at 4-bytes intervals and performs best when adjacent threads access data 4 bytes apart.
|
||||
hipSharedMemBankSizeEightByte ///< Shared mem is banked at 8-byte intervals and performs best when adjacent threads access data 4 bytes apart.
|
||||
} hipSharedMemConfig;
|
||||
|
||||
|
||||
|
||||
/**
|
||||
* Struct for data in 3D
|
||||
*
|
||||
*/
|
||||
typedef struct dim3 {
|
||||
uint32_t x; ///< x
|
||||
uint32_t y; ///< y
|
||||
uint32_t z; ///< z
|
||||
|
||||
dim3(uint32_t _x=1, uint32_t _y=1, uint32_t _z=1) : x(_x), y(_y), z(_z) {};
|
||||
} dim3;
|
||||
|
||||
|
||||
|
||||
/**
|
||||
* Memory copy types
|
||||
*
|
||||
*/
|
||||
enum hipMemcpyKind {
|
||||
hipMemcpyHostToHost = 0 ///< Host-to-Host Copy
|
||||
,hipMemcpyHostToDevice = 1 ///< Host-to-Device Copy
|
||||
,hipMemcpyDeviceToHost = 2 ///< Device-to-Host Copy
|
||||
,hipMemcpyDeviceToDevice =3 ///< Device-to-Device Copy
|
||||
,hipMemcpyDefault = 4, ///< Runtime will automatically determine copy-kind based on virtual addresses.
|
||||
} ;
|
||||
|
||||
|
||||
// Doxygen end group GlobalDefs
|
||||
/** @} */
|
||||
|
||||
|
||||
//-------------------------------------------------------------------------------------------------
|
||||
|
||||
|
||||
// The handle allows the async commands to use the stream even if the parent hipStream_t goes out-of-scope.
|
||||
typedef struct ihipStream_t * hipStream_t;
|
||||
|
||||
|
||||
/*
|
||||
* Opaque structure allows the true event (pointed at by the handle) to remain "live" even if the surrounding hipEvent_t goes out-of-scope.
|
||||
* This is handy for cases where the hipEvent_t goes out-of-scope but the true event is being written by some async queue or device */
|
||||
typedef struct hipEvent_t {
|
||||
struct ihipEvent_t *_handle;
|
||||
} hipEvent_t;
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
#ifdef __cplusplus
|
||||
} /* extern "C" */
|
||||
#endif
|
||||
|
||||
|
||||
|
||||
|
||||
//==================================================================================================
|
||||
#ifdef __cplusplus
|
||||
extern "C" {
|
||||
#endif
|
||||
|
||||
/**
|
||||
* @defgroup API HIP API
|
||||
* @{
|
||||
*
|
||||
* Defines the HIP API. See the individual sections for more information.
|
||||
*/
|
||||
|
||||
|
||||
|
||||
/**
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
* @defgroup Device Device Management
|
||||
* @{
|
||||
*/
|
||||
|
||||
/**
|
||||
* @brief Blocks until the default device has completed all preceding requested tasks.
|
||||
*
|
||||
* This function waits for all streams on the default device to complete execution, and then returns.
|
||||
*
|
||||
* @see hipSetDevice, hipDeviceReset
|
||||
*/
|
||||
hipError_t hipDeviceSynchronize(void);
|
||||
|
||||
|
||||
|
||||
/**
|
||||
* @brief Destroy all resources and reset all state on the default device in the current process.
|
||||
*
|
||||
* Explicity destroy all memory allocations, events, and queues associated with the default device in the current process.
|
||||
*
|
||||
* This function will reset the device immmediately, and then return after all resources have been freed.
|
||||
* The caller must ensure that the device is not being accessed by any other host threads from the active process when this function is called.
|
||||
*
|
||||
* @see hipDeviceSynchronize
|
||||
*/
|
||||
hipError_t hipDeviceReset(void) ;
|
||||
|
||||
|
||||
/**
|
||||
* @brief Set default device to be used for subsequent hip API calls from this thread.
|
||||
*
|
||||
* @param[in] device Valid device in range 0...hipGetDeviceCount().
|
||||
*
|
||||
* Sets @p device as the default device for the calling host thread. Valid device id's are 0... (hipGetDeviceCount()-1).
|
||||
*
|
||||
* Many HIP APIs implicitly use the "default device" :
|
||||
*
|
||||
* - Any device memory subsequently allocated from this host thread (using hipMalloc) will be allocated on device.
|
||||
* - Any streams or events created from this host thread will be associated with device.
|
||||
* - Any kernels launched from this host thread (using hipLaunchKernel) will be executed on device (unless a specific stream is specified,
|
||||
* in which case the device associated with that stream will be used).
|
||||
*
|
||||
* This function may be called from any host thread. Multiple host threads may use the same device.
|
||||
* This function does no synchronization with the previous or new device, and has very little runtime overhead.
|
||||
* Applications can use hipSetDevice to quickly switch the default device before making a HIP runtime call which uses the default device.
|
||||
*
|
||||
* The default device is stored in thread-local-storage for each thread.
|
||||
* Thread-pool implementations may inherit the default device of the previous thread. A good practice is to always call hipSetDevice
|
||||
* at the start of HIP coding sequency to establish a known standard device.
|
||||
*
|
||||
* @see hipGetDevice, hipGetDeviceCount
|
||||
*/
|
||||
hipError_t hipSetDevice(int device);
|
||||
|
||||
|
||||
/**
|
||||
* @brief Return the default device id for the calling host thread.
|
||||
*
|
||||
* @param [out] device *device is written with the default device
|
||||
*
|
||||
* HIP maintains an default device for each thread using thread-local-storage.
|
||||
* This device is used implicitly for HIP runtime APIs called by this thread.
|
||||
* hipGetDevice returns in * @p device the default device for the calling host thread.
|
||||
*
|
||||
* @see hipSetDevice, hipGetDevicesizeBytes
|
||||
*/
|
||||
hipError_t hipGetDevice(int *device);
|
||||
|
||||
|
||||
/**
|
||||
* @brief Return number of compute-capable devices.
|
||||
* @param [output] count Returns number of compute-capable devices.
|
||||
*
|
||||
* Returns in @p *count the number of devices that have ability to run compute commands. If there are no such devices, then @ref hipGetDeviceCount will return #hipErrorNoDevice.
|
||||
* If 1 or more devices can be found, then hipGetDeviceCount returns #hipSuccess.
|
||||
*/
|
||||
hipError_t hipGetDeviceCount(int *count);
|
||||
|
||||
|
||||
/**
|
||||
* @brief Returns device properties.
|
||||
*
|
||||
* @param [out] prop written with device properties
|
||||
* @param [in] device which device to query for information
|
||||
*
|
||||
* Populates hipDeviceGetProperties with information for the specified device.
|
||||
*/
|
||||
hipError_t hipDeviceGetProperties(hipDeviceProp_t* prop, int device);
|
||||
|
||||
|
||||
|
||||
//Cache partitioning functions:
|
||||
|
||||
/**
|
||||
* @brief Set L1/Shared cache partition.
|
||||
*
|
||||
* Note: AMD devices and recent NVIDIA GPUS do not support reconfigurable cache. This hint is ignored on those architectures.
|
||||
*
|
||||
*/
|
||||
hipError_t hipDeviceSetCacheConfig ( hipFuncCache cacheConfig );
|
||||
|
||||
|
||||
/**
|
||||
* @brief Set Cache configuration for a specific function
|
||||
*
|
||||
* Note: AMD devices and recent NVIDIA GPUS do not support reconfigurable cache. This hint is ignored on those architectures.
|
||||
*
|
||||
*/
|
||||
hipError_t hipDeviceGetCacheConfig ( hipFuncCache *cacheConfig );
|
||||
|
||||
|
||||
/**
|
||||
* @brief Set Cache configuration for a specific function
|
||||
*
|
||||
* Note: AMD devices and recent NVIDIA GPUS do not support reconfigurable cache. This hint is ignored on those architectures.
|
||||
*
|
||||
*/
|
||||
hipError_t hipFuncSetCacheConfig ( hipFuncCache config );
|
||||
|
||||
//---
|
||||
//Shared bank config functions:
|
||||
|
||||
/**
|
||||
* @brief Get Shared memory bank configuration.
|
||||
*
|
||||
* Note: AMD devices and recent NVIDIA GPUS do not support shared cache banking, and the hint is ignored on those architectures.
|
||||
*
|
||||
*/
|
||||
hipError_t hipDeviceGetSharedMemConfig ( hipSharedMemConfig * pConfig );
|
||||
|
||||
|
||||
/**
|
||||
* @brief Set Shared memory bank configuration.
|
||||
*
|
||||
* Note: AMD devices and recent NVIDIA GPUS do not support shared cache banking, and the hint is ignored on those architectures.
|
||||
*
|
||||
*/
|
||||
hipError_t hipDeviceSetSharedMemConfig ( hipSharedMemConfig config );
|
||||
|
||||
|
||||
// end doxygen Device
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
|
||||
/**
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
* @defgroup Error Error Handling
|
||||
* @{
|
||||
*/
|
||||
|
||||
/**
|
||||
* @brief Return last error returned by any HIP runtime API call and resets the stored error code to #hipSuccess
|
||||
*
|
||||
* Returns the last error that has been returned by any of the runtime calls in the same host thread,
|
||||
* and then resets the saved error to #hipSuccess.
|
||||
*
|
||||
*/
|
||||
hipError_t hipGetLastError(void);
|
||||
|
||||
|
||||
/**
|
||||
* @brief Return last error returned by any HIP runtime API call.
|
||||
*
|
||||
* @return #hipSuccess
|
||||
*
|
||||
* Returns the last error that has been returned by any of the runtime calls in the same host thread.
|
||||
* Unlike hipGetLastError, this function does not reset the saved error code.
|
||||
*
|
||||
*
|
||||
*
|
||||
*/
|
||||
hipError_t hipPeekAtLastError(void);
|
||||
|
||||
|
||||
|
||||
/**
|
||||
* @brief Return name of the specified error code in text form.
|
||||
*
|
||||
* @param hip_error Error code to convert to name.
|
||||
* @return const char pointer to the NULL-terminated error name
|
||||
*
|
||||
* @see hipGetErrorString, hipGetLastError, hipPeakAtLastError, hipError_t
|
||||
*/
|
||||
const char *hipGetErrorName(hipError_t hip_error);
|
||||
|
||||
|
||||
/**
|
||||
* @brief Return handy text string message to explain the error which occurred
|
||||
*
|
||||
* @param hip_error Error code to convert to string.
|
||||
* @return const char pointer to the NULL-terminated error string
|
||||
*
|
||||
* @warning : on HCC, this function returns the name of the error (same as hipGetErrorName)
|
||||
*
|
||||
* @see hipGetErrorName, hipGetLastError, hipPeakAtLastError, hipError_t
|
||||
*/
|
||||
const char *hipGetErrorString(hipError_t hip_error);
|
||||
|
||||
// end doxygen Error
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
|
||||
|
||||
/**
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
* @defgroup Stream Stream Management
|
||||
* @{
|
||||
*
|
||||
* The following Stream APIs are not (yet) supported in HIP:
|
||||
* - cudaStreamCreateWithPriority
|
||||
* - cudaStreamGetPriority
|
||||
*/
|
||||
|
||||
/**
|
||||
* @brief Create an asynchronous stream.
|
||||
*
|
||||
* @param[in, out] stream Pointer to new stream
|
||||
* @param[in ] flags to control stream creation.
|
||||
* @return #hipSuccess, #hipErrorInvalidValue
|
||||
*
|
||||
* Create a new asynchronous stream.
|
||||
* Flags controls behavior of the stream. See #hipStreamDefault, #hipStreamNonBlocking.
|
||||
* @error hipStream_t are under development - with current HIP use the NULL stream.
|
||||
*/
|
||||
|
||||
hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags);
|
||||
|
||||
|
||||
|
||||
/**
|
||||
* @brief Create an asynchronous stream.
|
||||
*
|
||||
* @param[in, out] stream Valid pointer to hipStream_t. This function writes the memory with the newly created stream.
|
||||
* @return #hipSuccess, #hipErrorInvalidValue
|
||||
*
|
||||
* Create a new asynchronous stream.
|
||||
*
|
||||
*/
|
||||
static inline hipError_t hipStreamCreate(hipStream_t *stream)
|
||||
{
|
||||
return hipStreamCreateWithFlags(stream, hipStreamDefault);
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* @brief Make the specified compute stream wait for an event
|
||||
*
|
||||
* @param[in] stream stream to make wait.
|
||||
* @param[in] event event to wait on
|
||||
* @param[in] flags control operation [must be 0]
|
||||
*
|
||||
* @return #hipSuccess, #hipErrorInvalidResourceHandle
|
||||
*
|
||||
* This function inserts a wait operation into the specified stream.
|
||||
* All future work submitted to @p stream will wait until @p event reports completion before beginning execution.
|
||||
* This function is host-asynchronous and the function may return before the wait has completed.
|
||||
*
|
||||
*
|
||||
*/
|
||||
hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags);
|
||||
|
||||
|
||||
/**
|
||||
* @brief Wait for all commands in stream to complete.
|
||||
*
|
||||
* If the null stream is specified, this command blocks until all
|
||||
*
|
||||
* This command honors the hipDeviceLaunchBlocking flag, which controls whether the wait is active or blocking.
|
||||
*
|
||||
* This command is host-synchronous : the host will block until the stream is empty.
|
||||
*
|
||||
* TODO
|
||||
*/
|
||||
hipError_t hipStreamSynchronize(hipStream_t stream);
|
||||
|
||||
|
||||
/**
|
||||
* @brief Destroys the specified stream.
|
||||
*
|
||||
* @param[in, out] stream Valid pointer to hipStream_t. This function writes the memory with the newly created stream.
|
||||
* @return #hipSuccess
|
||||
*
|
||||
* Destroys the specified stream.
|
||||
*
|
||||
* If commands are still executing on the specified stream, some may complete execution before the queue is deleted.
|
||||
*
|
||||
* The queue may be destroyed while some commands are still inflight, or may wait for all commands queued to the stream
|
||||
* before destroying it.
|
||||
*/
|
||||
hipError_t hipStreamDestroy(hipStream_t stream);
|
||||
|
||||
|
||||
/**
|
||||
* @brief Return flags associated with this stream.
|
||||
*
|
||||
* @param[in] stream
|
||||
* @param[in,out] flags
|
||||
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidResourceHandle
|
||||
*
|
||||
* Return flags associated with this stream in *@p flags.
|
||||
*
|
||||
* @see hipStreamCreateWithFlags
|
||||
*
|
||||
* @returns #hipSuccess
|
||||
*/
|
||||
hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags);
|
||||
|
||||
|
||||
// end doxygen Stream
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
|
||||
|
||||
|
||||
/**
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
* @defgroup Event Event Management
|
||||
* @{
|
||||
*/
|
||||
|
||||
/**
|
||||
* @brief Create an event with the specified flags
|
||||
*
|
||||
* @param[in,out] event Returns the newly created event.
|
||||
* @param[in] flags Flags to control event behavior. #hipEventDefault, #hipEventBlockingSync, #hipEventDisableTiming, #hipEventInterprocess
|
||||
*
|
||||
* @warning On HCC platform, #hipEventInterprocess is not supported.
|
||||
*
|
||||
* @returns #cudaSuccess
|
||||
*/
|
||||
hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags);
|
||||
|
||||
|
||||
/**
|
||||
* Create an event
|
||||
*
|
||||
* @param[in] event Creates an event
|
||||
*
|
||||
*/
|
||||
static inline hipError_t hipEventCreate(hipEvent_t* event)
|
||||
{
|
||||
return hipEventCreateWithFlags(event, 0);
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* @brief Record an event in the specified stream.
|
||||
*
|
||||
* @param[in] event event to record.
|
||||
* @param[in] stream stream in which to record event.
|
||||
* @returns #hipSuccess, #hipErrorInvalidResourceHandle
|
||||
*
|
||||
* hipEventQuery or hipEventSynchronize must be used to determine when the event
|
||||
* transitions from "recording" (after eventRecord is called) to "recorded"
|
||||
* (when timestamps are set, if requested).
|
||||
*
|
||||
* Events which are recorded in a non-NULL stream will transition to
|
||||
* from recording to "recorded" state when they reach the head of
|
||||
* the specified stream, after all previous
|
||||
* commands in that stream have completed executing.
|
||||
*
|
||||
* If hipEventRecord has been previously called aon event, then this call will overwrite any existing state in event.
|
||||
*
|
||||
* If this function is called on a an event that is currently being recorded, results are undefined - either
|
||||
* outstanding recording may save state into the event, and the order is not guaranteed. This shoul be avoided.
|
||||
*
|
||||
* @see hipEventElapsedTime
|
||||
*
|
||||
*/
|
||||
|
||||
hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream = NULL);
|
||||
|
||||
|
||||
/**
|
||||
* @brief Destroy the specified event.
|
||||
*
|
||||
* @param[in] event Event to destroy.
|
||||
* @return : #hipSuccess,
|
||||
*
|
||||
* Releases memory associated with the event. If the event is recording but has not completed recording when hipEventDestroy is called,
|
||||
* the function will return immediately and the completion_future resources will be released later, when the hipDevice is synchronized.
|
||||
*
|
||||
*/
|
||||
hipError_t hipEventDestroy(hipEvent_t event);
|
||||
|
||||
|
||||
/**
|
||||
* @brief: Wait for an event to complete.
|
||||
*
|
||||
* This function will block until the event is ready, waiting for all previous work in the stream specified when event was recorded with hipEventRecord.
|
||||
*
|
||||
* If hipEventRecord has not been called on @p event, this function returns immediately.
|
||||
*
|
||||
* TODO-hcc - This function needs to support hipEventBlockingSync parameter.
|
||||
*
|
||||
* @param[in] event Event on which to wait.
|
||||
* @return #hipSuccess, #hipErrorInvalidResourceHandle,
|
||||
*
|
||||
*/
|
||||
hipError_t hipEventSynchronize(hipEvent_t event);
|
||||
|
||||
|
||||
/**
|
||||
* @brief Return the elapsed time between two events.
|
||||
*
|
||||
* @param[out]] ms : Return time between start and stop in ms.
|
||||
* @param[in] start : Start event.
|
||||
* @param[in] stop : Stop event.
|
||||
* @return : #hipSuccess, #hipErrorInvalidResourceHandle, #hipErrorNotReady,
|
||||
*
|
||||
* Computes the elapsed time between two events. Time is computed in ms, with
|
||||
* a resolution of approximately 1 us.
|
||||
*
|
||||
* Events which are recorded in a NULL stream will block until all commands
|
||||
* on all other streams complete execution, and then record the timestamp.
|
||||
*
|
||||
* Events which are recorded in a non-NULL stream will record their timestamp
|
||||
* when they reach the head of the specified stream, after all previous
|
||||
* commands in that stream have completed executing. Thus the time that
|
||||
* the event recorded may be significantly after the host calls hipEventRecord.
|
||||
*
|
||||
* If hipEventRecord has not been called on either event, then #hipErrorInvalidResourceHandle is returned.
|
||||
* If hipEventRecord has been called on both events, but the timestamp has not yet been recorded on one or
|
||||
* both events (that is, hipEventQuery would return #hipErrorNotReady on at least one of the events), then
|
||||
* #hipErrorNotReady is returned.
|
||||
*/
|
||||
hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop);
|
||||
|
||||
|
||||
/**
|
||||
* @brief Query event status
|
||||
*
|
||||
* @param[in] event Event to query.
|
||||
* @returns #hipSuccess, hipEventNotReady
|
||||
*
|
||||
* Query the status of the specified event. This function will return #hipErrorNotReady if all commands
|
||||
* in the appropriate stream (specified to hipEventRecord) have completed. If that work has not completed,
|
||||
* or if hipEventRecord was not called on the event, then cudaSuccess is returned.
|
||||
*
|
||||
*
|
||||
*/
|
||||
hipError_t hipEventQuery(hipEvent_t event) ;
|
||||
|
||||
|
||||
// end doxygen Events
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
|
||||
|
||||
/**
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
* @defgroup Memory Memory Management
|
||||
* @{
|
||||
*
|
||||
* The following CUDA APIs are not currently supported:
|
||||
* - cudaMalloc3D
|
||||
* - cudaMalloc3DArray
|
||||
* - TODO - more 2D, 3D, array APIs here.
|
||||
*
|
||||
*
|
||||
*/
|
||||
|
||||
|
||||
|
||||
/**
|
||||
* Allocate memory on the default accelerator
|
||||
*
|
||||
* @param[out] ptr Pointer to the allocated memory
|
||||
* @param[in] size Requested memory size
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
hipError_t hipMalloc(void** ptr, size_t size) ;
|
||||
|
||||
|
||||
/**
|
||||
* Allocate pinned host memory
|
||||
*
|
||||
* @param[in] ptr Pointer to the allocated host pinned memory
|
||||
* @param[out] size Requested memory size
|
||||
* @return Error code
|
||||
*/
|
||||
hipError_t hipMallocHost(void** ptr, size_t size) ;
|
||||
|
||||
|
||||
|
||||
// TODO-doc (error codes)
|
||||
/**
|
||||
* Free memory allocated by the hcc hip memory allocation API.
|
||||
* This API performs an implicit hipDeviceSynchronize() call.
|
||||
*
|
||||
* @param[in] ptr Pointer to memory to be freed
|
||||
* @return Error code
|
||||
*/
|
||||
hipError_t hipFree(void* ptr);
|
||||
|
||||
|
||||
|
||||
// TODO-doc (error codes)
|
||||
/**
|
||||
* Free memory allocated by the hcc hip host memory allocation API
|
||||
*
|
||||
* @param[in] ptr Pointer to memory to be freed
|
||||
* @return Error code
|
||||
*/
|
||||
hipError_t hipFreeHost(void* ptr);
|
||||
|
||||
|
||||
|
||||
|
||||
// TODO-doc (error codes)
|
||||
/**
|
||||
* Copy data from src to dst. It supports memory from host to device,
|
||||
* device to host, device to device and host to host
|
||||
* The src and dst must not overlap.
|
||||
* If the
|
||||
*
|
||||
* This function is host-synchronous for most inputs.
|
||||
* It uses the default NULL stream and will synchronize with other blocking streams on the same device.
|
||||
*
|
||||
* @param[ being copy to
|
||||
* @param[in] src Data being copy from
|
||||
* @param[in] sizeBytes Data size in bytes
|
||||
* @param[in] copyType Memory copy type
|
||||
* @return Error code
|
||||
*/
|
||||
hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind);
|
||||
|
||||
|
||||
hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t sizeBytes, size_t offset, hipMemcpyKind kind);
|
||||
|
||||
|
||||
|
||||
// TODO-doc (error codes)
|
||||
/**
|
||||
* Copy data from src to dst asynchronously. It supports memory from host to device,
|
||||
* device to host, device to device and host to host.
|
||||
*
|
||||
* @param[out] dst Data being copy to
|
||||
* @param[in] src Data being copy from
|
||||
* @param[in] sizeBytes Data size in bytes
|
||||
* @param[in] accelerator_view Accelerator view which the copy is being enqueued
|
||||
* @return Error code
|
||||
*/
|
||||
hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream=0);
|
||||
|
||||
|
||||
// TODO-doc
|
||||
/*
|
||||
* This function is host-asynchronous and may return before the memset operation completes.
|
||||
* Same as hipMemsetAsync with null stream.
|
||||
*
|
||||
* */
|
||||
hipError_t hipMemset(void* dst, int value, size_t sizeBytes );
|
||||
hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t = 0 );
|
||||
|
||||
|
||||
/*
|
||||
* @brief Query memory info. Return snapshot of free memory, and total allocatable memory on the device.
|
||||
*
|
||||
* Returns in *free a snapshot of the current free memory o
|
||||
**/
|
||||
hipError_t hipMemGetInfo (size_t * free, size_t * total) ;
|
||||
|
||||
// doxygen end Memory
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
|
||||
|
||||
/**
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
* @defgroup PeerToPeer Device Memory Access
|
||||
* @{
|
||||
*
|
||||
*/
|
||||
|
||||
/**
|
||||
* @brief Determine if a device can access a peer's memory.
|
||||
*
|
||||
* @param [out] canAccessPeer returns true if specified devices are peers.
|
||||
* @param [in] device
|
||||
* @param [in] peerDevice
|
||||
*
|
||||
* Returns "1" in @p canAccessPeer if the specified @p device is capable
|
||||
* of directly accessing memory phyically located on peerDevice , or "0" if not.
|
||||
*/
|
||||
hipError_t hipDeviceCanAccessPeer ( int* canAccessPeer, int device, int peerDevice );
|
||||
|
||||
// TODO-DOC
|
||||
hipError_t hipDeviceDisablePeerAccess ( int peerDevice );
|
||||
|
||||
// TODO-DOC
|
||||
hipError_t hipDeviceEnablePeerAccess ( int peerDevice, unsigned int flags );
|
||||
|
||||
// TODO-DOC
|
||||
hipError_t hipMemcpyPeer ( void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes );
|
||||
|
||||
// TODO-DOC
|
||||
hipError_t hipMemcpyPeerAsync ( void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes, hipStream_t stream=0 );
|
||||
// doxygen end PeerToPeer
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
|
||||
/**
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
* @defgroup Version Management
|
||||
* @{
|
||||
*
|
||||
*/
|
||||
|
||||
/**
|
||||
* @brief Returns the approximate HIP driver versin.
|
||||
*
|
||||
* @warning The HIP feature set does not correpond to an exact CUDA SDK driver revision.
|
||||
* This function always set *driverVersion to 4 as an approximation though HIP supports
|
||||
* some features which were introduced in later CUDA SDK revisions.
|
||||
* HIP apps code should not rely on the driver revision number here and should
|
||||
* use arch feature flags to test device capabiliies or conditional compilation.
|
||||
*
|
||||
*/
|
||||
hipError_t hipDriverGetVersion(int *driverVersion) ;
|
||||
|
||||
|
||||
|
||||
// doxygen end Version Management
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
|
||||
/**
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
* @defgroup Profiler Control
|
||||
* @{
|
||||
*
|
||||
*
|
||||
* The cudaProfilerInitialize API format for "configFile" is not supported.
|
||||
*
|
||||
* On AMD platforms, hipProfilerStart and hipProfilerStop require installation of AMD's GPU
|
||||
* perf counter API and defining GPU_PERF
|
||||
*/
|
||||
|
||||
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
|
||||
|
||||
|
||||
#ifdef __cplusplus
|
||||
} /* extern "c" */
|
||||
#endif
|
||||
|
||||
|
||||
/**
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
* @defgroup HCC_Specific HCC-Specific Accessors
|
||||
* @{
|
||||
*
|
||||
* The following calls are only supported when compiler HIP with HCC.
|
||||
* To produce portable code, use of these calls must be guarded #ifdef checks:
|
||||
* @code
|
||||
* #ifdef __HCC__
|
||||
* hc::accelerator acc;
|
||||
hipError_t err = hipHccGetAccelerator(deviceId, &acc)
|
||||
* #endif
|
||||
* @endcode
|
||||
*
|
||||
*/
|
||||
|
||||
#ifdef __HCC__
|
||||
#include <hc.hpp>
|
||||
/**
|
||||
* @brief Return hc::acclerator associated with the specified deviceId
|
||||
*/
|
||||
hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator *acc);
|
||||
|
||||
/**
|
||||
* @brief Return hc::acclerator_view associated with the specified stream
|
||||
*/
|
||||
hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **av);
|
||||
#endif
|
||||
|
||||
|
||||
// end-group HCC_Specific
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
|
||||
|
||||
// doxygen end HIP API
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
@@ -0,0 +1,179 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
#pragma once
|
||||
/**
|
||||
* @file hip_kalmar_texture.h
|
||||
* @brief HIP C++ Texture API for hcc compiler
|
||||
*/
|
||||
|
||||
#include <limits.h>
|
||||
|
||||
#include <hip_runtime.h>
|
||||
|
||||
//----
|
||||
//Texture - TODO - likely need to move this to a separate file only included with kernel compilation.
|
||||
#define hipTextureType1D 1
|
||||
|
||||
typedef struct hipChannelFormatDesc {
|
||||
// TODO - this has 4-5 well-defined fields, we could just copy...
|
||||
int _dummy;
|
||||
} hipChannelFormatDesc;
|
||||
|
||||
typedef enum hipTextureReadMode
|
||||
{
|
||||
hipReadModeElementType, ///< Read texture as specified element type
|
||||
//! @warning cudaReadModeNormalizedFloat is not supported.
|
||||
} hipTextureReadMode;
|
||||
|
||||
typedef enum hipTextureFilterMode
|
||||
{
|
||||
hipFilterModePoint, ///< Point filter mode.
|
||||
//! @warning cudaFilterModeLinear is not supported.
|
||||
} hipTextureFilterMode;
|
||||
|
||||
struct textureReference {
|
||||
hipTextureFilterMode filterMode;
|
||||
bool normalized;
|
||||
hipChannelFormatDesc channelDesc;
|
||||
};
|
||||
|
||||
template <class T, int texType=hipTextureType1D, enum hipTextureReadMode=hipReadModeElementType>
|
||||
struct texture : public textureReference {
|
||||
|
||||
const T * _dataPtr; // pointer to underlying data.
|
||||
|
||||
//texture() : filterMode(hipFilterModePoint), normalized(false), _dataPtr(NULL) {};
|
||||
};
|
||||
|
||||
|
||||
|
||||
|
||||
#define tex1Dfetch(_tex, _addr) (_tex._dataPtr[_addr])
|
||||
|
||||
|
||||
|
||||
|
||||
/**
|
||||
* @addtogroup API HIP API
|
||||
* @{
|
||||
*
|
||||
* Defines the HIP API. See the individual sections for more information.
|
||||
*/
|
||||
|
||||
// These are C++ APIs - maybe belong in separate file.
|
||||
/**
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
* @defgroup Texture Texture Reference Management
|
||||
* @{
|
||||
*
|
||||
*
|
||||
* @warning The HIP texture API implements a small subset of full texture API. Known limitations include:
|
||||
* - Only point sampling is supported.
|
||||
* - Only C++ APIs are provided.
|
||||
* - Many APIs and modes are not implemented.
|
||||
*
|
||||
* The HIP texture support is intended to allow use of texture cache on hardware where this is beneficial.
|
||||
*
|
||||
* The following CUDA APIs are not currently supported:
|
||||
* - cudaBindTexture2D
|
||||
* - cudaBindTextureToArray
|
||||
* - cudaBindTextureToMipmappedArray
|
||||
* - cudaGetChannelDesc
|
||||
* - cudaGetTextureReference
|
||||
*
|
||||
*/
|
||||
|
||||
// C API:
|
||||
#if 0
|
||||
hipChannelFormatDesc hipBindTexture(size_t *offset, struct textureReference *tex, const void *devPtr, const struct hipChannelFormatDesc *desc, size_t size=UINT_MAX)
|
||||
{
|
||||
tex->_dataPtr = devPtr;
|
||||
}
|
||||
#endif
|
||||
|
||||
/*
|
||||
* @brief hipChannelFormatDesc
|
||||
**/
|
||||
// TODO
|
||||
template <class T>
|
||||
hipChannelFormatDesc hipCreateChannelDesc()
|
||||
{
|
||||
hipChannelFormatDesc desc;
|
||||
return desc;
|
||||
}
|
||||
|
||||
/*
|
||||
* @brief hipBindTexture
|
||||
**/
|
||||
// TODO-doc
|
||||
template <class T, int dim, enum hipTextureReadMode readMode>
|
||||
hipError_t hipBindTexture(size_t *offset,
|
||||
struct texture<T, dim, readMode> &tex,
|
||||
const void *devPtr,
|
||||
const struct hipChannelFormatDesc *desc,
|
||||
size_t size=UINT_MAX)
|
||||
{
|
||||
tex._dataPtr = static_cast<const T*>(devPtr);
|
||||
|
||||
return hipSuccess;
|
||||
}
|
||||
|
||||
|
||||
/*
|
||||
* @brief hipBindTexture
|
||||
**/
|
||||
// TODO-doc
|
||||
template <class T, int dim, enum hipTextureReadMode readMode>
|
||||
hipError_t hipBindTexture(size_t *offset,
|
||||
struct texture<T, dim, readMode> &tex,
|
||||
const void *devPtr,
|
||||
size_t size=UINT_MAX)
|
||||
{
|
||||
return hipBindTexture(offset, tex, devPtr, &tex.channelDesc, size);
|
||||
}
|
||||
|
||||
|
||||
/*
|
||||
* @brief hipUnbindTexture
|
||||
**/
|
||||
// TODO-doc
|
||||
template <class T, int dim, enum hipTextureReadMode readMode>
|
||||
hipError_t hipUnbindTexture(struct texture<T, dim, readMode> *tex)
|
||||
{
|
||||
tex->_dataPtr = NULL;
|
||||
|
||||
return hipSuccess;
|
||||
}
|
||||
|
||||
|
||||
|
||||
// doxygen end Texture
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
|
||||
// End doxygen API:
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
@@ -0,0 +1,187 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
#if defined (__HCC__) && (__hcc_workweek__ < 16032)
|
||||
#error("This version of HIP requires a newer version of HCC.");
|
||||
#endif
|
||||
|
||||
#include <hc_short_vector.hpp>
|
||||
|
||||
//-- Signed
|
||||
// Define char vector types
|
||||
typedef hc::short_vector::char1 char1;
|
||||
typedef hc::short_vector::char2 char2;
|
||||
typedef hc::short_vector::char3 char3;
|
||||
typedef hc::short_vector::char4 char4;
|
||||
|
||||
// Define short vector types
|
||||
typedef hc::short_vector::short1 short1;
|
||||
typedef hc::short_vector::short2 short2;
|
||||
typedef hc::short_vector::short3 short3;
|
||||
typedef hc::short_vector::short4 short4;
|
||||
|
||||
// Define int vector types
|
||||
typedef hc::short_vector::int1 int1;
|
||||
typedef hc::short_vector::int2 int2;
|
||||
typedef hc::short_vector::int3 int3;
|
||||
typedef hc::short_vector::int4 int4;
|
||||
|
||||
// Define long vector types
|
||||
typedef hc::short_vector::long1 long1;
|
||||
typedef hc::short_vector::long2 long2;
|
||||
typedef hc::short_vector::long3 long3;
|
||||
typedef hc::short_vector::long4 long4;
|
||||
|
||||
// Define longlong vector types
|
||||
typedef hc::short_vector::longlong1 longlong1;
|
||||
typedef hc::short_vector::longlong2 longlong2;
|
||||
typedef hc::short_vector::longlong3 longlong3;
|
||||
typedef hc::short_vector::longlong4 longlong4;
|
||||
|
||||
|
||||
//-- Unsigned
|
||||
// Define uchar vector types
|
||||
typedef hc::short_vector::uchar1 uchar1;
|
||||
typedef hc::short_vector::uchar2 uchar2;
|
||||
typedef hc::short_vector::uchar3 uchar3;
|
||||
typedef hc::short_vector::uchar4 uchar4;
|
||||
|
||||
// Define ushort vector types
|
||||
typedef hc::short_vector::ushort1 ushort1;
|
||||
typedef hc::short_vector::ushort2 ushort2;
|
||||
typedef hc::short_vector::ushort3 ushort3;
|
||||
typedef hc::short_vector::ushort4 ushort4;
|
||||
|
||||
// Define uint vector types
|
||||
typedef hc::short_vector::uint1 uint1;
|
||||
typedef hc::short_vector::uint2 uint2;
|
||||
typedef hc::short_vector::uint3 uint3;
|
||||
typedef hc::short_vector::uint4 uint4;
|
||||
|
||||
// Define ulong vector types
|
||||
typedef hc::short_vector::ulong1 ulong1;
|
||||
typedef hc::short_vector::ulong2 ulong2;
|
||||
typedef hc::short_vector::ulong3 ulong3;
|
||||
typedef hc::short_vector::ulong4 ulong4;
|
||||
|
||||
// Define ulonglong vector types
|
||||
typedef hc::short_vector::ulonglong1 ulonglong1;
|
||||
typedef hc::short_vector::ulonglong2 ulonglong2;
|
||||
typedef hc::short_vector::ulonglong3 ulonglong3;
|
||||
typedef hc::short_vector::ulonglong4 ulonglong4;
|
||||
|
||||
|
||||
//-- Floating point
|
||||
// Define float vector types
|
||||
typedef hc::short_vector::float1 float1;
|
||||
typedef hc::short_vector::float2 float2;
|
||||
typedef hc::short_vector::float3 float3;
|
||||
typedef hc::short_vector::float4 float4;
|
||||
|
||||
// Define double vector types
|
||||
typedef hc::short_vector::double1 double1;
|
||||
typedef hc::short_vector::double2 double2;
|
||||
typedef hc::short_vector::double3 double3;
|
||||
typedef hc::short_vector::double4 double4;
|
||||
|
||||
|
||||
///---
|
||||
// Inline functions for creating vector types from basic types
|
||||
#define ONE_COMPONENT_ACCESS(T, VT) inline VT make_ ##VT (T x) { VT t; t.x = x; return t; };
|
||||
#define TWO_COMPONENT_ACCESS(T, VT) inline VT make_ ##VT (T x, T y) { VT t; t.x=x; t.y=y; return t; };
|
||||
#define THREE_COMPONENT_ACCESS(T, VT) inline VT make_ ##VT (T x, T y, T z) { VT t; t.x=x; t.y=y; t.z=z; return t; };
|
||||
#define FOUR_COMPONENT_ACCESS(T, VT) inline VT make_ ##VT (T x, T y, T z, T w) { VT t; t.x=x; t.y=y; t.z=z; t.w=w; return t; };
|
||||
|
||||
|
||||
//signed:
|
||||
ONE_COMPONENT_ACCESS (signed char, char1);
|
||||
TWO_COMPONENT_ACCESS (signed char, char2);
|
||||
THREE_COMPONENT_ACCESS(signed char, char3);
|
||||
FOUR_COMPONENT_ACCESS (signed char, char4);
|
||||
|
||||
ONE_COMPONENT_ACCESS (short, short1);
|
||||
TWO_COMPONENT_ACCESS (short, short2);
|
||||
THREE_COMPONENT_ACCESS(short, short3);
|
||||
FOUR_COMPONENT_ACCESS (short, short4);
|
||||
|
||||
ONE_COMPONENT_ACCESS (int, int1);
|
||||
TWO_COMPONENT_ACCESS (int, int2);
|
||||
THREE_COMPONENT_ACCESS(int, int3);
|
||||
FOUR_COMPONENT_ACCESS (int, int4);
|
||||
|
||||
ONE_COMPONENT_ACCESS (long int, long1);
|
||||
TWO_COMPONENT_ACCESS (long int, long2);
|
||||
THREE_COMPONENT_ACCESS(long int, long3);
|
||||
FOUR_COMPONENT_ACCESS (long int, long4);
|
||||
|
||||
ONE_COMPONENT_ACCESS (long long int, ulong1);
|
||||
TWO_COMPONENT_ACCESS (long long int, ulong2);
|
||||
THREE_COMPONENT_ACCESS(long long int, ulong3);
|
||||
FOUR_COMPONENT_ACCESS (long long int, ulong4);
|
||||
|
||||
ONE_COMPONENT_ACCESS (long long int, longlong1);
|
||||
TWO_COMPONENT_ACCESS (long long int, longlong2);
|
||||
THREE_COMPONENT_ACCESS(long long int, longlong3);
|
||||
FOUR_COMPONENT_ACCESS (long long int, longlong4);
|
||||
|
||||
|
||||
// unsigned:
|
||||
ONE_COMPONENT_ACCESS (unsigned char, uchar1);
|
||||
TWO_COMPONENT_ACCESS (unsigned char, uchar2);
|
||||
THREE_COMPONENT_ACCESS(unsigned char, uchar3);
|
||||
FOUR_COMPONENT_ACCESS (unsigned char, uchar4);
|
||||
|
||||
ONE_COMPONENT_ACCESS (unsigned short, ushort1);
|
||||
TWO_COMPONENT_ACCESS (unsigned short, ushort2);
|
||||
THREE_COMPONENT_ACCESS(unsigned short, ushort3);
|
||||
FOUR_COMPONENT_ACCESS (unsigned short, ushort4);
|
||||
|
||||
ONE_COMPONENT_ACCESS (unsigned int, uint1);
|
||||
TWO_COMPONENT_ACCESS (unsigned int, uint2);
|
||||
THREE_COMPONENT_ACCESS(unsigned int, uint3);
|
||||
FOUR_COMPONENT_ACCESS (unsigned int, uint4);
|
||||
|
||||
ONE_COMPONENT_ACCESS (unsigned long int, ulong1);
|
||||
TWO_COMPONENT_ACCESS (unsigned long int, ulong2);
|
||||
THREE_COMPONENT_ACCESS(unsigned long int, ulong3);
|
||||
FOUR_COMPONENT_ACCESS (unsigned long int, ulong4);
|
||||
|
||||
ONE_COMPONENT_ACCESS (unsigned long long int, ulong1);
|
||||
TWO_COMPONENT_ACCESS (unsigned long long int, ulong2);
|
||||
THREE_COMPONENT_ACCESS(unsigned long long int, ulong3);
|
||||
FOUR_COMPONENT_ACCESS (unsigned long long int, ulong4);
|
||||
|
||||
ONE_COMPONENT_ACCESS (unsigned long long int, ulonglong1);
|
||||
TWO_COMPONENT_ACCESS (unsigned long long int, ulonglong2);
|
||||
THREE_COMPONENT_ACCESS(unsigned long long int, ulonglong3);
|
||||
FOUR_COMPONENT_ACCESS (unsigned long long int, ulonglong4);
|
||||
|
||||
|
||||
//Floating point
|
||||
ONE_COMPONENT_ACCESS (float, float1);
|
||||
TWO_COMPONENT_ACCESS (float, float2);
|
||||
THREE_COMPONENT_ACCESS(float, float3);
|
||||
FOUR_COMPONENT_ACCESS (float, float4);
|
||||
|
||||
ONE_COMPONENT_ACCESS (double, double1);
|
||||
TWO_COMPONENT_ACCESS (double, double2);
|
||||
THREE_COMPONENT_ACCESS(double, double3);
|
||||
FOUR_COMPONENT_ACCESS (double, double4);
|
||||
@@ -0,0 +1,63 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
#ifdef __HCC__
|
||||
/**
|
||||
* Function and kernel markers
|
||||
*/
|
||||
#define __host__ __attribute__((cpu))
|
||||
#define __device__ __attribute__((hc))
|
||||
|
||||
#ifndef DISABLE_GRID_LAUNCH
|
||||
#define __global__ __attribute__((hc_grid_launch))
|
||||
#else
|
||||
#define __global__
|
||||
#endif
|
||||
|
||||
#define __noinline__ __attribute__((noinline))
|
||||
#define __forceinline__ __attribute__((always_inline))
|
||||
|
||||
|
||||
|
||||
/*
|
||||
* Variable Type Qualifiers:
|
||||
*/
|
||||
// _restrict is supported by the compiler
|
||||
#define __shared__ tile_static
|
||||
#define __constant__ __attribute__((address_space(2)))
|
||||
|
||||
#else
|
||||
// Non-HCC compiler
|
||||
/**
|
||||
* Function and kernel markers
|
||||
*/
|
||||
#define __host__
|
||||
#define __device__
|
||||
|
||||
#define __global__
|
||||
|
||||
#define __noinline__
|
||||
#define __forceinline__
|
||||
|
||||
#define __shared__
|
||||
#define __constant__
|
||||
|
||||
#endif
|
||||
@@ -0,0 +1,85 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
|
||||
// Disable use of grid_launch feature in HCC compiler.
|
||||
//#define DISABLE_GRID_LAUNCH
|
||||
|
||||
// Common code included at start of every hip file.
|
||||
// Auto enable __HIP_PLATFORM_HCC__ if compiling with HCC
|
||||
// Other compiler (GCC,ICC,etc) need to set one of these macros explicitly
|
||||
#if defined(__HCC__)
|
||||
#define __HIP_PLATFORM_HCC__
|
||||
#define __HIPCC__
|
||||
|
||||
# if defined __HCC_ACCELERATOR__
|
||||
# define __HIP_DEVICE_COMPILE__ 1
|
||||
# endif
|
||||
#endif
|
||||
|
||||
// Auto enable __HIP_PLATFORM_NVCC__ if compiling with NVCC
|
||||
#if defined(__NVCC__)
|
||||
#define __HIP_PLATFORM_NVCC__
|
||||
# ifdef __CUDACC__
|
||||
# define __HIPCC__
|
||||
# endif
|
||||
|
||||
# ifdef __CUDA_ARCH__
|
||||
# define __HIP_DEVICE_COMPILE__ 1
|
||||
# endif
|
||||
|
||||
#endif
|
||||
|
||||
|
||||
|
||||
|
||||
#if __HIP_DEVICE_COMPILE__ == 0
|
||||
// 32-bit Atomics
|
||||
#define __HIP_ARCH_HAS_GLOBAL_INT32_ATOMICS__ (0)
|
||||
#define __HIP_ARCH_HAS_GLOBAL_FLOAT_ATOMIC_EXCH__ (0)
|
||||
#define __HIP_ARCH_HAS_SHARED_INT32_ATOMICS__ (0)
|
||||
#define __HIP_ARCH_HAS_SHARED_FLOAT_ATOMIC_EXCH__ (0)
|
||||
#define __HIP_ARCH_HAS_FLOAT_ATOMIC_ADD__ (0)
|
||||
|
||||
// 64-bit Atomics
|
||||
#define __HIP_ARCH_HAS_GLOBAL_INT64_ATOMICS__ (0)
|
||||
#define __HIP_ARCH_HAS_SHARED_INT64_ATOMICS__ (0)
|
||||
|
||||
// Doubles
|
||||
#define __HIP_ARCH_HAS_DOUBLES__ (0)
|
||||
|
||||
// Warp cross-lane operations
|
||||
#define __HIP_ARCH_HAS_WARP_VOTE__ (0)
|
||||
#define __HIP_ARCH_HAS_WARP_BALLOT__ (0)
|
||||
#define __HIP_ARCH_HAS_WARP_SHUFFLE__ (0)
|
||||
#define __HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ (0)
|
||||
|
||||
// Sync
|
||||
#define __HIP_ARCH_HAS_THREAD_FENCE_SYSTEM__ (0)
|
||||
#define __HIP_ARCH_HAS_SYNC_THREAD_EXT__ (0)
|
||||
|
||||
// Misc
|
||||
#define __HIP_ARCH_HAS_SURFACE_FUNCS__ (0)
|
||||
#define __HIP_ARCH_HAS_3DGRID__ (0)
|
||||
#define __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ (0)
|
||||
#endif
|
||||
@@ -0,0 +1,59 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
//! HIP = Heterogeneous-compute Interface for Portability
|
||||
//!
|
||||
//! Define a extremely thin runtime layer that allows source code to be compiled unmodified
|
||||
//! through either AMD HCC or NVCC. Key features tend to be in the spirit
|
||||
//! and terminology of CUDA, but with a portable path to other accelerators as well:
|
||||
//
|
||||
//! Both paths support rich C++ features including classes, templates, lambdas, etc.
|
||||
//! Runtime API is C
|
||||
//! Memory management is based on pure pointers and resembles malloc/free/copy.
|
||||
//
|
||||
//! hip_runtime.h : includes everything in hip_api.h, plus math builtins and kernel launch macros.
|
||||
//! hip_runtime_api.h : Defines HIP API. This is a C header file and does not use any C++ features.
|
||||
|
||||
#pragma once
|
||||
|
||||
// Some standard header files, these are included by hc.hpp and so want to make them avail on both
|
||||
// paths to provide a consistent include env and avoid "missing symbol" errors that only appears
|
||||
// on NVCC path:
|
||||
#include <stdint.h>
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <assert.h>
|
||||
|
||||
#include <hip_common.h>
|
||||
|
||||
#if defined(__HIP_PLATFORM_HCC__) and not defined (__HIP_PLATFORM_NVCC__)
|
||||
#include <hcc_detail/hip_runtime.h>
|
||||
#elif defined(__HIP_PLATFORM_NVCC__) and not defined (__HIP_PLATFORM_HCC__)
|
||||
#include <nvcc_detail/hip_runtime.h>
|
||||
#else
|
||||
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
|
||||
#endif
|
||||
|
||||
|
||||
#include <hip_runtime_api.h>
|
||||
#include <hip_vector_types.h>
|
||||
|
||||
@@ -0,0 +1,162 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
/**
|
||||
* @file hip_runtime_api.h
|
||||
*
|
||||
* Defines the API signatures for HIP runtime.
|
||||
* This file can be compiled with a standard compiler.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
|
||||
|
||||
#include <string.h> // for getDeviceProp
|
||||
#include <hip_common.h>
|
||||
|
||||
typedef struct {
|
||||
// 32-bit Atomics:
|
||||
unsigned hasGlobalInt32Atomics : 1; ///< 32-bit integer atomics for global memory
|
||||
unsigned hasGlobalFloatAtomicExch : 1; ///< 32-bit float atomic exch for global memory
|
||||
unsigned hasSharedInt32Atomics : 1; ///< 32-bit integer atomics for shared memory
|
||||
unsigned hasSharedFloatAtomicExch : 1; ///< 32-bit float atomic exch for shared memory
|
||||
unsigned hasFloatAtomicAdd : 1; ///< 32-bit float atomic add in global and shared memory
|
||||
|
||||
// 64-bit Atomics:
|
||||
unsigned hasGlobalInt64Atomics : 1; ///< 64-bit integer atomics for global memory
|
||||
unsigned hasSharedInt64Atomics : 1; ///< 64-bit integer atomics for shared memory
|
||||
|
||||
// Doubles
|
||||
unsigned hasDoubles : 1; ///< double-precision floating point.
|
||||
|
||||
// Warp cross-lane operations:
|
||||
unsigned hasWarpVote : 1; ///< warp vote instructions (__any, __all)
|
||||
unsigned hasWarpBallot : 1; ///< warp ballot instructions (__ballot)
|
||||
unsigned hasWarpShuffle : 1; ///< warp shuffle operations. (__shfl_*)
|
||||
unsigned hasFunnelShift : 1; ///< funnel two words into one, with shift&mask caps
|
||||
|
||||
// Sync
|
||||
unsigned hasThreadFenceSystem : 1; ///< __threadfence_system
|
||||
unsigned hasSyncThreadsExt : 1; ///< __syncthreads_count, syncthreads_and, syncthreads_or
|
||||
|
||||
// Misc
|
||||
unsigned hasSurfaceFuncs : 1; ///< Surface functions
|
||||
unsigned has3dGrid : 1; ///< Grid and group dims are 3D (rather than 2D)
|
||||
unsigned hasDynamicParallelism : 1; ///< Dynamic parallellism
|
||||
} hipDeviceArch_t;
|
||||
|
||||
|
||||
//---
|
||||
// Common headers for both NVCC and HCC paths:
|
||||
|
||||
/**
|
||||
* hipDeviceProp
|
||||
*
|
||||
*/
|
||||
typedef struct hipDeviceProp_t {
|
||||
char name[256]; ///< Device name
|
||||
size_t totalGlobalMem; ///< Size of global memory region (in bytes)
|
||||
size_t sharedMemPerBlock; ///< Size of shared memory region (in bytes)
|
||||
int regsPerBlock ; ///< registers per block
|
||||
int warpSize ; ///< warp size
|
||||
int maxThreadsPerBlock; ///< max work items per work group or workgroup max size
|
||||
int maxThreadsDim[3]; ///< max number of threads in each dimension (XYZ) of a block
|
||||
int maxGridSize[3]; ///< max grid dimensions (XYZ)
|
||||
int clockRate ; ///< max clock frequency of the multiProcessors, in khz.
|
||||
|
||||
size_t totalConstMem; ///< Size of shared memory region (in bytes)
|
||||
int major ; ///< Major compute capability. On HCC, this is an approximation and features may differ from CUDA CC. See the arch feature flags for portable ways to query feature caps.
|
||||
int minor; ///< Minor compute capability. On HCC, this is an approximation and features may differ from CUDA CC. See the arch feature flags for portable ways to query feature caps.
|
||||
int multiProcessorCount; ///< number of multi-processors (compute units)
|
||||
int l2CacheSize; ///< L2 cache size
|
||||
int maxThreadsPerMultiProcessor; ///< Maximum resident threads per multi-processor
|
||||
int computeMode; ///< Compute mode
|
||||
|
||||
int clockInstructionRate ; ///< Frequency in khz of the timer used by the device-side "clock*" instructions. New for HIP.
|
||||
|
||||
hipDeviceArch_t arch; ///< Architectural feature flags. New for HIP.
|
||||
} hipDeviceProp_t;
|
||||
|
||||
|
||||
// hack to get these to show up in Doxygen:
|
||||
/**
|
||||
* @defgroup GlobalDefs Global enum and defines
|
||||
* @{
|
||||
*
|
||||
*/
|
||||
|
||||
|
||||
/*
|
||||
* @brief hipError_t
|
||||
* @enum
|
||||
* @ingroup Enumerations
|
||||
*/
|
||||
typedef enum hipError_t {
|
||||
hipSuccess = 0 ///< Successful completion.
|
||||
,hipErrorMemoryAllocation ///< Memory allocation error.
|
||||
,hipErrorMemoryFree ///< Memory free error.
|
||||
,hipErrorUnknownSymbol ///< Unknown symbol
|
||||
,hipErrorOutOfResources ///< Out of resources error
|
||||
,hipErrorInvalidValue ///< One or more of the paramters passed to the API call is NULL or not in an acceptable range.
|
||||
,hipErrorInvalidResourceHandle ///< Resource handle (hipEvent_t or hipStream_t) invalid.
|
||||
,hipErrorInvalidDevice ///< DeviceID must be in range 0...#compute-devices.
|
||||
,hipErrorNoDevice ///< Call to cudaGetDeviceCount returned 0 devices
|
||||
,hipErrorNotReady ///< indicates that asynchronous operations enqueued earlier are not ready. This is not actually an error, but is used to distinguish from hipSuccess (which indicates completion). APIs that return this error include hipEventQuery and hipStreamQuery.
|
||||
|
||||
,hipErrorUnknown ///< Unknown error
|
||||
,hipErrorTbd ///< Marker that more error codes are needed.
|
||||
} hipError_t;
|
||||
|
||||
|
||||
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
#if defined(__HIP_PLATFORM_HCC__) and not defined (__HIP_PLATFORM_NVCC__)
|
||||
#include "hcc_detail/hip_runtime_api.h"
|
||||
#elif defined(__HIP_PLATFORM_NVCC__) and not defined (__HIP_PLATFORM_HCC__)
|
||||
#include "nvcc_detail/hip_runtime_api.h"
|
||||
#else
|
||||
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
|
||||
#endif
|
||||
|
||||
|
||||
/**
|
||||
* @brief: C++ wrapper for hipMalloc
|
||||
*
|
||||
* Perform automatic type conversion to eliminate need for excessive typecasting (ie void**)
|
||||
*
|
||||
* @see hipMalloc
|
||||
*/
|
||||
#ifdef __cplusplus
|
||||
template<class T>
|
||||
static inline hipError_t hipMalloc ( T** devPtr, size_t size)
|
||||
{
|
||||
return hipMalloc((void**)devPtr, size);
|
||||
}
|
||||
|
||||
template<class T>
|
||||
static inline hipError_t hipMallocHost ( T** ptr, size_t size)
|
||||
{
|
||||
return hipMallocHost((void**)ptr, size);
|
||||
}
|
||||
#endif
|
||||
@@ -0,0 +1,35 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
//! hip_vector_types.h : Defines the HIP vector types.
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <hip_common.h>
|
||||
|
||||
|
||||
#if defined(__HIP_PLATFORM_HCC__) and not defined (__HIP_PLATFORM_NVCC__)
|
||||
#include <hcc_detail/hip_vector_types.h>
|
||||
#elif defined(__HIP_PLATFORM_NVCC__) and not defined (__HIP_PLATFORM_HCC__)
|
||||
#include <vector_types.h>
|
||||
#else
|
||||
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
|
||||
#endif
|
||||
@@ -0,0 +1,105 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
#pragma once
|
||||
|
||||
#include <cuda_runtime.h>
|
||||
|
||||
#include <hip_runtime_api.h>
|
||||
|
||||
#define HIP_KERNEL_NAME(...) __VA_ARGS__
|
||||
|
||||
typedef int hipLaunchParm ;
|
||||
|
||||
#define hipLaunchKernel(kernelName, numblocks, numthreads, memperblock, streamId, ...) \
|
||||
do {\
|
||||
kernelName<<<numblocks,numthreads,memperblock,streamId>>>(0, __VA_ARGS__);\
|
||||
} while(0)
|
||||
|
||||
|
||||
#define KERNELBEGIN
|
||||
#define KERNELEND
|
||||
|
||||
#define hipReadModeElementType cudaReadModeElementType
|
||||
|
||||
#ifdef __CUDA_ARCH__
|
||||
|
||||
|
||||
// 32-bit Atomics:
|
||||
#define __HIP_ARCH_HAS_GLOBAL_INT32_ATOMICS__ (__CUDA_ARCH__ >= 110)
|
||||
#define __HIP_ARCH_HAS_GLOBAL_FLOAT_ATOMIC_EXCH__ (__CUDA_ARCH__ >= 110)
|
||||
#define __HIP_ARCH_HAS_SHARED_INT32_ATOMICS__ (__CUDA_ARCH__ >= 120)
|
||||
#define __HIP_ARCH_HAS_SHARED_FLOAT_ATOMIC_EXCH__ (__CUDA_ARCH__ >= 120)
|
||||
#define __HIP_ARCH_HAS_FLOAT_ATOMIC_ADD__
|
||||
|
||||
// 64-bit Atomics:
|
||||
#define __HIP_ARCH_HAS_GLOBAL_INT64_ATOMICS__ (__CUDA_ARCH__ >= 200)
|
||||
#define __HIP_ARCH_HAS_SHARED_INT64_ATOMICS__ (__CUDA_ARCH__ >= 120)
|
||||
|
||||
// Doubles
|
||||
#define __HIP_ARCH_HAS_DOUBLES__ (__CUDA_ARCH__ >= 120)
|
||||
|
||||
//warp cross-lane operations:
|
||||
#define __HIP_ARCH_HAS_WARP_VOTE__ (__CUDA_ARCH__ >= 120)
|
||||
#define __HIP_ARCH_HAS_WARP_BALLOT__ (__CUDA_ARCH__ >= 200)
|
||||
#define __HIP_ARCH_HAS_WARP_SHUFFLE__ (__CUDA_ARCH__ >= 300)
|
||||
#define __HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ (__CUDA_ARCH__ >= 350)
|
||||
|
||||
//sync
|
||||
#define __HIP_ARCH_HAS_THREAD_FENCE_SYSTEM__ (__CUDA_ARCH__ >= 200)
|
||||
#define __HIP_ARCH_HAS_SYNC_THREAD_EXT__ (__CUDA_ARCH__ >= 200)
|
||||
|
||||
// misc
|
||||
#define __HIP_ARCH_HAS_SURFACE_FUNCS__ (__CUDA_ARCH__ >= 200)
|
||||
#define __HIP_ARCH_HAS_3DGRID__ (__CUDA_ARCH__ >= 200)
|
||||
#define __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ (__CUDA_ARCH__ >= 350)
|
||||
|
||||
|
||||
#else
|
||||
|
||||
#define __HIP_DEVICE_COMPILE__ 0
|
||||
|
||||
#endif
|
||||
|
||||
#ifdef __CUDACC__
|
||||
|
||||
|
||||
|
||||
|
||||
#define hipThreadIdx_x threadIdx.x
|
||||
#define hipThreadIdx_y threadIdx.y
|
||||
#define hipThreadIdx_z threadIdx.z
|
||||
|
||||
#define hipBlockIdx_x blockIdx.x
|
||||
#define hipBlockIdx_y blockIdx.y
|
||||
#define hipBlockIdx_z blockIdx.z
|
||||
|
||||
#define hipBlockDim_x blockDim.x
|
||||
#define hipBlockDim_y blockDim.y
|
||||
#define hipBlockDim_z blockDim.z
|
||||
|
||||
#define hipGridDim_x gridDim.x
|
||||
#define hipGridDim_y gridDim.y
|
||||
#define hipGridDim_z gridDim.z
|
||||
|
||||
#endif
|
||||
|
||||
|
||||
@@ -0,0 +1,338 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
#pragma once
|
||||
|
||||
#include <cuda_runtime_api.h>
|
||||
|
||||
|
||||
#ifdef __cplusplus
|
||||
extern "C" {
|
||||
#endif
|
||||
|
||||
//TODO -move to include/hip_runtime_api.h as a common implementation.
|
||||
/**
|
||||
* Memory copy types
|
||||
*
|
||||
*/
|
||||
typedef enum hipMemcpyKind {
|
||||
hipMemcpyHostToHost
|
||||
,hipMemcpyHostToDevice
|
||||
,hipMemcpyDeviceToHost
|
||||
,hipMemcpyDeviceToDevice
|
||||
,hipMemcpyDefault
|
||||
} hipMemcpyKind ;
|
||||
|
||||
// hipErrorNoDevice.
|
||||
|
||||
/*typedef enum hipTextureFilterMode
|
||||
{
|
||||
hipFilterModePoint = cudaFilterModePoint, ///< Point filter mode.
|
||||
//! @warning cudaFilterModeLinear is not supported.
|
||||
} hipTextureFilterMode;*/
|
||||
#define hipFilterModePoint cudaFilterModePoint
|
||||
|
||||
|
||||
typedef cudaEvent_t hipEvent_t;
|
||||
typedef cudaStream_t hipStream_t;
|
||||
//typedef cudaChannelFormatDesc hipChannelFormatDesc;
|
||||
#define hipChannelFormatDesc cudaChannelFormatDesc
|
||||
|
||||
inline static hipError_t hipCUDAErrorTohipError(cudaError_t cuError) {
|
||||
switch(cuError) {
|
||||
case cudaSuccess:
|
||||
return hipSuccess;
|
||||
case cudaErrorMemoryAllocation:
|
||||
return hipErrorMemoryAllocation;
|
||||
case cudaErrorInvalidDevicePointer:
|
||||
case cudaErrorInitializationError:
|
||||
return hipErrorMemoryFree;
|
||||
default:
|
||||
return hipErrorUnknown;
|
||||
}
|
||||
}
|
||||
// TODO match the error enum names of hip and cuda
|
||||
inline static cudaError_t hipErrorToCudaError(hipError_t hError) {
|
||||
switch(hError) {
|
||||
case hipSuccess:
|
||||
return cudaSuccess;
|
||||
case hipErrorMemoryAllocation:
|
||||
return cudaErrorMemoryAllocation;
|
||||
case hipErrorMemoryFree:
|
||||
return cudaErrorInitializationError;
|
||||
default:
|
||||
return cudaErrorUnknown;
|
||||
}
|
||||
}
|
||||
|
||||
inline static cudaMemcpyKind hipMemcpyKindToCudaMemcpyKind(hipMemcpyKind kind) {
|
||||
switch(kind) {
|
||||
case hipMemcpyHostToHost:
|
||||
return cudaMemcpyHostToHost;
|
||||
case hipMemcpyHostToDevice:
|
||||
return cudaMemcpyHostToDevice;
|
||||
case hipMemcpyDeviceToHost:
|
||||
return cudaMemcpyDeviceToHost;
|
||||
default:
|
||||
return cudaMemcpyDefault;
|
||||
}
|
||||
}
|
||||
|
||||
inline static hipError_t hipDeviceReset() {
|
||||
return hipCUDAErrorTohipError(cudaDeviceReset());
|
||||
}
|
||||
|
||||
inline static hipError_t hipGetLastError() {
|
||||
return hipCUDAErrorTohipError(cudaGetLastError());
|
||||
}
|
||||
|
||||
inline static hipError_t hipMalloc(void** ptr, size_t size) {
|
||||
return hipCUDAErrorTohipError(cudaMalloc(ptr, size));
|
||||
}
|
||||
|
||||
inline static hipError_t hipFree(void* ptr) {
|
||||
return hipCUDAErrorTohipError(cudaFree(ptr));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMallocHost(void** ptr, size_t size) {
|
||||
return hipCUDAErrorTohipError(cudaMallocHost(ptr, size));
|
||||
}
|
||||
inline static hipError_t hipFreeHost(void* ptr) {
|
||||
return hipCUDAErrorTohipError(cudaFreeHost(ptr));
|
||||
}
|
||||
inline static hipError_t hipSetDevice(int device) {
|
||||
return hipCUDAErrorTohipError(cudaSetDevice(device));
|
||||
}
|
||||
inline static hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind copyKind) {
|
||||
return hipCUDAErrorTohipError(cudaMemcpy(dst, src, sizeBytes, hipMemcpyKindToCudaMemcpyKind(copyKind)));
|
||||
}
|
||||
|
||||
|
||||
inline static hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind copyKind, hipStream_t stream=0) {
|
||||
return hipCUDAErrorTohipError(cudaMemcpyAsync(dst, src, sizeBytes, hipMemcpyKindToCudaMemcpyKind(copyKind), stream));
|
||||
}
|
||||
|
||||
|
||||
inline static hipError_t hipMemcpyToSymbol(const char * symbolName, const void* src, size_t sizeBytes, size_t offset = 0, hipMemcpyKind copyType = hipMemcpyHostToDevice) {
|
||||
return hipCUDAErrorTohipError(cudaMemcpyToSymbol(symbolName, src, sizeBytes, offset, hipMemcpyKindToCudaMemcpyKind(copyType)));
|
||||
}
|
||||
inline static hipError_t hipDeviceSynchronize() {
|
||||
return hipCUDAErrorTohipError(cudaDeviceSynchronize());
|
||||
}
|
||||
|
||||
inline static const char* hipGetErrorString(hipError_t error){
|
||||
return cudaGetErrorString( hipErrorToCudaError(error) );
|
||||
}
|
||||
|
||||
inline static hipError_t hipGetDeviceCount(int * count){
|
||||
return hipCUDAErrorTohipError(cudaGetDeviceCount(count));
|
||||
}
|
||||
|
||||
inline static hipError_t hipGetDevice(int * device){
|
||||
return hipCUDAErrorTohipError(cudaGetDevice(device));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMemset(void* devPtr,int value, size_t count) {
|
||||
return hipCUDAErrorTohipError(cudaMemset(devPtr, value, count));
|
||||
}
|
||||
|
||||
inline static hipError_t hipDeviceGetProperties(hipDeviceProp_t *p_prop, int device)
|
||||
{
|
||||
cudaDeviceProp cdprop;
|
||||
cudaError_t cerror;
|
||||
cerror = cudaGetDeviceProperties(&cdprop,device);
|
||||
strcpy(p_prop->name,cdprop.name);
|
||||
p_prop->totalGlobalMem = cdprop.totalGlobalMem ;
|
||||
p_prop->sharedMemPerBlock = cdprop.sharedMemPerBlock;
|
||||
p_prop->regsPerBlock = cdprop.regsPerBlock;
|
||||
p_prop->warpSize = cdprop.warpSize ;
|
||||
for (int i=0 ; i<3; i++) {
|
||||
p_prop->maxThreadsDim[i] = cdprop.maxThreadsDim[i];
|
||||
p_prop->maxGridSize[i] = cdprop.maxGridSize[i];
|
||||
}
|
||||
p_prop->maxThreadsPerBlock = cdprop.maxThreadsPerBlock ;
|
||||
p_prop->clockRate = cdprop.clockRate;
|
||||
p_prop->totalConstMem = cdprop.totalConstMem ;
|
||||
p_prop->major = cdprop.major ;
|
||||
p_prop->minor = cdprop. minor ;
|
||||
p_prop->multiProcessorCount = cdprop.multiProcessorCount ;
|
||||
p_prop->l2CacheSize = cdprop.l2CacheSize ;
|
||||
p_prop->maxThreadsPerMultiProcessor = cdprop.maxThreadsPerMultiProcessor ;
|
||||
p_prop->computeMode = cdprop.computeMode ;
|
||||
|
||||
// Same as clock-rate:
|
||||
p_prop->clockInstructionRate = cdprop.clockRate;
|
||||
|
||||
int ccVers = p_prop->major*100 + p_prop->minor * 10;
|
||||
|
||||
p_prop->arch.hasGlobalInt32Atomics = (ccVers >= 110);
|
||||
p_prop->arch.hasGlobalFloatAtomicExch = (ccVers >= 110);
|
||||
p_prop->arch.hasSharedInt32Atomics = (ccVers >= 120);
|
||||
p_prop->arch.hasSharedFloatAtomicExch = (ccVers >= 120);
|
||||
|
||||
p_prop->arch.hasFloatAtomicAdd = (ccVers >= 200);
|
||||
|
||||
p_prop->arch.hasGlobalInt64Atomics = (ccVers >= 120);
|
||||
p_prop->arch.hasSharedInt64Atomics = (ccVers >= 110);
|
||||
|
||||
p_prop->arch.hasDoubles = (ccVers >= 130);
|
||||
|
||||
p_prop->arch.hasWarpVote = (ccVers >= 120);
|
||||
p_prop->arch.hasWarpBallot = (ccVers >= 200);
|
||||
p_prop->arch.hasWarpShuffle = (ccVers >= 300);
|
||||
p_prop->arch.hasFunnelShift = (ccVers >= 350);
|
||||
|
||||
p_prop->arch.hasThreadFenceSystem = (ccVers >= 200);
|
||||
p_prop->arch.hasSyncThreadsExt = (ccVers >= 200);
|
||||
|
||||
p_prop->arch.hasSurfaceFuncs = (ccVers >= 200);
|
||||
p_prop->arch.has3dGrid = (ccVers >= 200);
|
||||
p_prop->arch.hasDynamicParallelism = (ccVers >= 350);
|
||||
|
||||
|
||||
return hipCUDAErrorTohipError(cerror);
|
||||
}
|
||||
|
||||
inline static hipError_t hipMemGetInfo( size_t* free, size_t* total)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaMemGetInfo(free,total));
|
||||
}
|
||||
|
||||
inline static hipError_t hipEventCreate( hipEvent_t* event)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaEventCreate(event));
|
||||
}
|
||||
|
||||
inline static hipError_t hipEventRecord( hipEvent_t event, hipStream_t stream = NULL)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaEventRecord(event,stream));
|
||||
}
|
||||
|
||||
inline static hipError_t hipEventSynchronize( hipEvent_t event)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaEventSynchronize(event));
|
||||
}
|
||||
|
||||
inline static hipError_t hipEventElapsedTime( float *ms, hipEvent_t start, hipEvent_t stop)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaEventElapsedTime(ms,start,stop));
|
||||
}
|
||||
|
||||
inline static hipError_t hipEventDestroy( hipEvent_t event)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaEventDestroy(event));
|
||||
}
|
||||
|
||||
|
||||
inline static hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaStreamCreateWithFlags(stream, flags));
|
||||
}
|
||||
|
||||
|
||||
inline static hipError_t hipStreamCreate(hipStream_t *stream)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaStreamCreate(stream));
|
||||
}
|
||||
|
||||
|
||||
inline static hipError_t hipStreamDestroy(hipStream_t stream)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaStreamDestroy(stream));
|
||||
}
|
||||
|
||||
|
||||
inline static hipError_t hipDriverGetVersion(int *driverVersion)
|
||||
{
|
||||
cudaError_t err = cudaDriverGetVersion(driverVersion);
|
||||
|
||||
// Override driver version to match version reported on HCC side.
|
||||
*driverVersion = 4;
|
||||
|
||||
return hipCUDAErrorTohipError(err);
|
||||
}
|
||||
|
||||
|
||||
inline static hipError_t hipDeviceCanAccessPeer ( int* canAccessPeer, int device, int peerDevice )
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaDeviceCanAccessPeer(canAccessPeer, device, peerDevice));
|
||||
}
|
||||
|
||||
inline static hipError_t hipDeviceDisablePeerAccess ( int peerDevice )
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaDeviceDisablePeerAccess ( peerDevice ));
|
||||
};
|
||||
|
||||
inline static hipError_t hipDeviceEnablePeerAccess ( int peerDevice, unsigned int flags )
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaDeviceEnablePeerAccess ( peerDevice, flags ));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMemcpyPeer ( void* dst, int dstDevice, const void* src, int srcDevice, size_t count )
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaMemcpyPeer ( dst, dstDevice, src, srcDevice, count ));
|
||||
};
|
||||
|
||||
inline static hipError_t hipMemcpyPeerAsync ( void* dst, int dstDevice, const void* src, int srcDevice, size_t count, hipStream_t stream=0 )
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaMemcpyPeerAsync ( dst, dstDevice, src, srcDevice, count, stream ));
|
||||
};
|
||||
|
||||
|
||||
|
||||
|
||||
#ifdef __cplusplus
|
||||
}
|
||||
#endif
|
||||
|
||||
#ifdef __CUDACC__
|
||||
|
||||
template <class T, int dim, enum cudaTextureReadMode readMode>
|
||||
inline static hipError_t hipBindTexture(size_t *offset,
|
||||
const struct texture<T, dim, readMode> &tex,
|
||||
const void *devPtr,
|
||||
size_t size=UINT_MAX)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, size));
|
||||
}
|
||||
|
||||
template <class T, int dim, enum cudaTextureReadMode readMode>
|
||||
inline static hipError_t hipBindTexture(size_t *offset,
|
||||
struct texture<T, dim, readMode> *tex,
|
||||
const void *devPtr,
|
||||
const struct hipChannelFormatDesc *desc,
|
||||
size_t size=UINT_MAX)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size));
|
||||
}
|
||||
|
||||
template <class T, int dim, enum cudaTextureReadMode readMode>
|
||||
inline static hipError_t hipUnbindTexture(struct texture<T, dim, readMode> *tex)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaUnbindTexture(tex));
|
||||
}
|
||||
|
||||
template <class T>
|
||||
inline static hipChannelFormatDesc hipCreateChannelDesc()
|
||||
{
|
||||
return cudaCreateChannelDesc<T>();
|
||||
}
|
||||
#endif
|
||||
Reference in New Issue
Block a user