AI/rocm-systems
2
0
Vork 0
Code Kwesties Pull-aanvragen Actions Paketten Projecten Publicaties Wiki Activiteit

C guarding the apis

Bladeren bron 
[ROCm/hip commit: 496da0aabd]
This commit is contained in:
Aditya Atluri
2016-03-30 06:08:50 -05:00
bovenliggende 3683e2a3bd
commit afd8f2bcfc
10 gewijzigde bestanden met toevoegingen van 64 en 97 verwijderingen
Download Patch-bestand Download Diff-bestand Expand all files Collapse all files
projects/hip/include/hcc_detail/hip_runtime.h
+6 -70
Bestand weergeven
@@ -30,8 +30,8 @@ THE SOFTWARE.
// Top part of file can be compiled with any compiler
#include <cstring>
#include <cmath>
//#include <cstring>
//#include <cmath>
#include <string.h>
#include <stddef.h>
@@ -43,7 +43,9 @@ THE SOFTWARE.
//---
// Remainder of this file only compiles with HCC
#ifdef __HCC__
//#if __cplusplus
#include <hc.hpp>
//#endif
#include <grid_launch.h>
extern int HIP_TRACE_API;
@@ -54,12 +56,12 @@ extern int HIP_TRACE_API;
#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__)
#if defined (__KALMAR_ACCELERATOR__) && !defined (__HCC_ACCELERATOR__)
#define __HCC_ACCELERATOR__ __KALMAR_ACCELERATOR__
#endif
// Feature tests:
#if defined(__HCC_ACCELERATOR__) and (__HCC_ACCELERATOR__ != 0)
#if defined(__HCC_ACCELERATOR__) && (__HCC_ACCELERATOR__ != 0)
// Device compile and not host compile:
//TODO-HCC enable __HIP_ARCH_HAS_ATOMICS__ when HCC supports these.
@@ -472,36 +474,8 @@ __device__ inline float __dsqrt_rz(double x) {return hc::fast_math::sqrt(x); };
#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__
hipStream_t ihipPreLaunchKernel(hipStream_t stream, hc::accelerator_view **av);
void ihipPostLaunchKernel(hipStream_t stream, hc::completion_future &cf);
@@ -565,44 +539,6 @@ do {\
#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__
projects/hip/include/hcc_detail/hip_runtime_api.h
+23 -11
Bestand weergeven
@@ -114,8 +114,9 @@ typedef struct dim3 {
uint32_t x; ///< x
uint32_t y; ///< y
uint32_t z; ///< z
#ifdef __cplusplus
dim3(uint32_t _x=1, uint32_t _y=1, uint32_t _z=1) : x(_x), y(_y), z(_z) {};
#endif
} dim3;
@@ -124,13 +125,13 @@ typedef struct dim3 {
* Memory copy types
*
*/
enum hipMemcpyKind {
typedef 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.
} ;
} hipMemcpyKind;
@@ -577,9 +578,11 @@ hipError_t hipEventCreate(hipEvent_t* event);
* @see hipEventElapsedTime
*
*/
#ifdef __cplusplus
hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream = NULL);
#else
hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream);
#endif
/**
* @brief Destroy the specified event.
@@ -677,8 +680,7 @@ hipError_t hipEventQuery(hipEvent_t event) ;
/**
* @brief Return attributes for the specified pointer
*/
hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr) ;
hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr);
/**
* @brief Allocate memory on the default accelerator
@@ -845,8 +847,11 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t siz
* @param[in] accelerator_view Accelerator view which the copy is being enqueued
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree, #hipErrorUnknown
*/
#if __cplusplus
hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream=0);
#else
hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream);
#endif
/**
* @brief Copy data from src to dst asynchronously.
@@ -876,8 +881,11 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes );
* @param[in] stream - Stream identifier
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree
*/
#if __cplusplus
hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t = 0 );
#else
hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream);
#endif
/**
* @brief Query memory info.
@@ -964,7 +972,11 @@ hipError_t hipMemcpyPeer ( void* dst, int dstDevice, const void* src, int srcD
*
* Returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidDevice
*/
#if __cplusplus
hipError_t hipMemcpyPeerAsync ( void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes, hipStream_t stream=0 );
#else
hipError_t hipMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes, hipStream_t stream);
#endif
// doxygen end PeerToPeer
/**
* @}
@@ -1041,7 +1053,7 @@ hipError_t hipDriverGetVersion(int *driverVersion) ;
* @endcode
*
*/
#if __cplusplus
#ifdef __HCC__
#include <hc.hpp>
/**
@@ -1054,7 +1066,7 @@ hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator *acc);
*/
hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **av);
#endif
#endif
// end-group HCC_Specific
/**
projects/hip/include/hcc_detail/hip_texture.h
+2 -2
Bestand weergeven
@@ -55,7 +55,7 @@ struct textureReference {
bool normalized;
hipChannelFormatDesc channelDesc;
};
#if __cplusplus
template <class T, int texType=hipTextureType1D, enum hipTextureReadMode=hipReadModeElementType>
struct texture : public textureReference {
@@ -63,7 +63,7 @@ struct texture : public textureReference {
//texture() : filterMode(hipFilterModePoint), normalized(false), _dataPtr(NULL) {};
};
#endif
projects/hip/include/hip_common.h
+2 -2
Bestand weergeven
@@ -32,7 +32,7 @@ THE SOFTWARE.
#define __HIP_PLATFORM_HCC__
#define __HIPCC__
#if defined(__HCC_ACCELERATOR__) and (__HCC_ACCELERATOR__ != 0)
#if defined(__HCC_ACCELERATOR__) && (__HCC_ACCELERATOR__ != 0)
#define __HIP_DEVICE_COMPILE__ 1
#else
#define __HIP_DEVICE_COMPILE__ 0
@@ -46,7 +46,7 @@ THE SOFTWARE.
# define __HIPCC__
# endif
#if defined(__CUDA_ARCH__) and (__CUDA_ARCH__ != 0)
#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ != 0)
#define __HIP_DEVICE_COMPILE__ 1
#else
#define __HIP_DEVICE_COMPILE__ 0
projects/hip/include/hip_runtime.h
+2 -2
Bestand weergeven
@@ -50,9 +50,9 @@ THE SOFTWARE.
#include <hip_common.h>
#if defined(__HIP_PLATFORM_HCC__) and not defined (__HIP_PLATFORM_NVCC__)
#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
#include <hcc_detail/hip_runtime.h>
#elif defined(__HIP_PLATFORM_NVCC__) and not defined (__HIP_PLATFORM_HCC__)
#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__)
#include <nvcc_detail/hip_runtime.h>
#else
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
projects/hip/include/hip_runtime_api.h
+2 -2
Bestand weergeven
@@ -197,9 +197,9 @@ typedef enum hipDeviceAttribute_t {
* @}
*/
#if defined(__HIP_PLATFORM_HCC__) and not defined (__HIP_PLATFORM_NVCC__)
#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
#include "hcc_detail/hip_runtime_api.h"
#elif defined(__HIP_PLATFORM_NVCC__) and not defined (__HIP_PLATFORM_HCC__)
#elif defined(__HIP_PLATFORM_NVCC__) && !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__");
projects/hip/include/hip_vector_types.h
+2 -2
Bestand weergeven
@@ -26,9 +26,9 @@ THE SOFTWARE.
#include <hip_common.h>
#if defined(__HIP_PLATFORM_HCC__) and not defined (__HIP_PLATFORM_NVCC__)
#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
#include <hcc_detail/hip_vector_types.h>
#elif defined(__HIP_PLATFORM_NVCC__) and not defined (__HIP_PLATFORM_HCC__)
#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__)
#include <vector_types.h>
#else
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
projects/hip/tests/src/hipC.c
+2 -2
Bestand weergeven
@@ -3,7 +3,7 @@
#define ITER 1<<20
#define SIZE 1024*1024*sizeof(int)
/*
__global__ void Iter(hipLaunchParm lp, int *Ad){
int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
if(tx == 0){
@@ -12,7 +12,7 @@ __global__ void Iter(hipLaunchParm lp, int *Ad){
}
}
}
*/
int main(){
int A=0, *Ad;
hipMalloc((void**)&Ad, SIZE);
projects/hip/tests/src/hipCKernel.c
+19
Bestand weergeven
@@ -0,0 +1,19 @@
#include<hip_runtime.h>
#include<hip_runtime_api.h>
__global__ void Kernel(hipLaunchParm lp, float *Ad){
int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
}
int main(){
dim3 dimBlock;
dim3 dimGrid;
dimGrid.x = 1;
dimGrid.y = 1;
dimGrid.z = 1;
dimBlock.x = 1;
dimBlock.y = 1;
dimBlock.z = 1;
float *A;
hipLaunchKernel(HIP_KERNEL_NAME(Kernel), dimGrid, dimBlock, 0, 0, A);
}
projects/hip/tests/src/hipLanguageExtensions.cpp
+4 -4
Bestand weergeven
@@ -53,14 +53,14 @@ __device__ __host__ float PlusOne(float x)
__global__ void MyKernel (const hipLaunchParm lp, const float *a, const float *b, float *c, unsigned N)
{
KERNELBEGIN;
//KERNELBEGIN;
unsigned gid = hipThreadIdx_x;
if (gid < N) {
c[gid] = a[gid] + PlusOne(b[gid]);
}
KERNELEND;
//KERNELEND;
}
@@ -82,7 +82,7 @@ vectorADD(const hipLaunchParm lp,
T *C_d,
size_t N)
{
KERNELBEGIN;
// KERNELBEGIN;
int ws = warpSize;
@@ -119,7 +119,7 @@ vectorADD(const hipLaunchParm lp,
C_d[i] = A_d[i] + B_d[i];
}
KERNELEND;
// KERNELEND;
}