From afd8f2bcfcf69b10a8a0e59939f409722696d10a Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Wed, 30 Mar 2016 06:08:50 -0500 Subject: [PATCH] C guarding the apis [ROCm/hip commit: 496da0aabd31593a6cfe234be785282a13fc9298] --- projects/hip/include/hcc_detail/hip_runtime.h | 76 ++----------------- .../hip/include/hcc_detail/hip_runtime_api.h | 34 ++++++--- projects/hip/include/hcc_detail/hip_texture.h | 4 +- projects/hip/include/hip_common.h | 4 +- projects/hip/include/hip_runtime.h | 4 +- projects/hip/include/hip_runtime_api.h | 4 +- projects/hip/include/hip_vector_types.h | 4 +- projects/hip/tests/src/hipC.c | 4 +- projects/hip/tests/src/hipCKernel.c | 19 +++++ .../hip/tests/src/hipLanguageExtensions.cpp | 8 +- 10 files changed, 64 insertions(+), 97 deletions(-) create mode 100644 projects/hip/tests/src/hipCKernel.c diff --git a/projects/hip/include/hcc_detail/hip_runtime.h b/projects/hip/include/hcc_detail/hip_runtime.h index a691ae8913..7328ff4362 100644 --- a/projects/hip/include/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hcc_detail/hip_runtime.h @@ -30,8 +30,8 @@ THE SOFTWARE. // Top part of file can be compiled with any compiler -#include -#include +//#include +//#include #include #include @@ -43,7 +43,9 @@ THE SOFTWARE. //--- // Remainder of this file only compiles with HCC #ifdef __HCC__ +//#if __cplusplus #include +//#endif #include extern int HIP_TRACE_API; @@ -54,12 +56,12 @@ extern int HIP_TRACE_API; #include #include // 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__ diff --git a/projects/hip/include/hcc_detail/hip_runtime_api.h b/projects/hip/include/hcc_detail/hip_runtime_api.h index 544c16cd83..8ae92c3625 100644 --- a/projects/hip/include/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hcc_detail/hip_runtime_api.h @@ -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 /** @@ -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 /** diff --git a/projects/hip/include/hcc_detail/hip_texture.h b/projects/hip/include/hcc_detail/hip_texture.h index 595d110274..1a20f1960b 100644 --- a/projects/hip/include/hcc_detail/hip_texture.h +++ b/projects/hip/include/hcc_detail/hip_texture.h @@ -55,7 +55,7 @@ struct textureReference { bool normalized; hipChannelFormatDesc channelDesc; }; - +#if __cplusplus template struct texture : public textureReference { @@ -63,7 +63,7 @@ struct texture : public textureReference { //texture() : filterMode(hipFilterModePoint), normalized(false), _dataPtr(NULL) {}; }; - +#endif diff --git a/projects/hip/include/hip_common.h b/projects/hip/include/hip_common.h index d8be2bb928..919105fef5 100644 --- a/projects/hip/include/hip_common.h +++ b/projects/hip/include/hip_common.h @@ -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 diff --git a/projects/hip/include/hip_runtime.h b/projects/hip/include/hip_runtime.h index de2a49d975..61a02cb3ea 100644 --- a/projects/hip/include/hip_runtime.h +++ b/projects/hip/include/hip_runtime.h @@ -50,9 +50,9 @@ THE SOFTWARE. #include -#if defined(__HIP_PLATFORM_HCC__) and not defined (__HIP_PLATFORM_NVCC__) +#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__) #include -#elif defined(__HIP_PLATFORM_NVCC__) and not defined (__HIP_PLATFORM_HCC__) +#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__) #include #else #error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__"); diff --git a/projects/hip/include/hip_runtime_api.h b/projects/hip/include/hip_runtime_api.h index 3a0a4b399a..bef881d066 100644 --- a/projects/hip/include/hip_runtime_api.h +++ b/projects/hip/include/hip_runtime_api.h @@ -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__"); diff --git a/projects/hip/include/hip_vector_types.h b/projects/hip/include/hip_vector_types.h index 731aa7193f..61d7f6c7d9 100644 --- a/projects/hip/include/hip_vector_types.h +++ b/projects/hip/include/hip_vector_types.h @@ -26,9 +26,9 @@ THE SOFTWARE. #include -#if defined(__HIP_PLATFORM_HCC__) and not defined (__HIP_PLATFORM_NVCC__) +#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__) #include -#elif defined(__HIP_PLATFORM_NVCC__) and not defined (__HIP_PLATFORM_HCC__) +#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__) #include #else #error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__"); diff --git a/projects/hip/tests/src/hipC.c b/projects/hip/tests/src/hipC.c index 73526ecfe0..7f95682c39 100644 --- a/projects/hip/tests/src/hipC.c +++ b/projects/hip/tests/src/hipC.c @@ -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); diff --git a/projects/hip/tests/src/hipCKernel.c b/projects/hip/tests/src/hipCKernel.c new file mode 100644 index 0000000000..15f12fcadd --- /dev/null +++ b/projects/hip/tests/src/hipCKernel.c @@ -0,0 +1,19 @@ +#include +#include + +__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); +} diff --git a/projects/hip/tests/src/hipLanguageExtensions.cpp b/projects/hip/tests/src/hipLanguageExtensions.cpp index 361189e258..9c484ca64f 100644 --- a/projects/hip/tests/src/hipLanguageExtensions.cpp +++ b/projects/hip/tests/src/hipLanguageExtensions.cpp @@ -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; }