From 3b7a8fab5b525babe4ab008744edf1f4e2e7437e Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Mon, 25 Sep 2017 17:55:49 +0530 Subject: [PATCH 1/8] Added Texture 2D object and reference APIs tests --- tests/src/texture/hipTextureObj2D.cpp | 107 ++++++++++++++++++++++++++ tests/src/texture/hipTextureRef2D.cpp | 100 ++++++++++++++++++++++++ 2 files changed, 207 insertions(+) create mode 100644 tests/src/texture/hipTextureObj2D.cpp create mode 100644 tests/src/texture/hipTextureRef2D.cpp diff --git a/tests/src/texture/hipTextureObj2D.cpp b/tests/src/texture/hipTextureObj2D.cpp new file mode 100644 index 0000000000..bea21d0658 --- /dev/null +++ b/tests/src/texture/hipTextureObj2D.cpp @@ -0,0 +1,107 @@ +#include +#include +#include + +#include +#include "test_common.h" + +bool testResult = true; + +__global__ void tex2DKernel(float* outputData, + hipTextureObject_t textureObject, + int width, + int height) +{ + int x = hipBlockIdx_x*hipBlockDim_x + hipThreadIdx_x; + int y = hipBlockIdx_y*hipBlockDim_y + hipThreadIdx_y; + outputData[y*width + x] = tex2D(textureObject, x, y); +} + +void runTest(int argc, char **argv); + +int main(int argc, char **argv) +{ + runTest(argc, argv); + + if(testResult) { + passed(); + } else { + exit(EXIT_FAILURE); + } + +} + +void runTest(int argc, char **argv) +{ + unsigned int width = 256; + unsigned int height = 256; + unsigned int size = width * height * sizeof(float); + float* hData = (float*) malloc(size); + memset(hData, 0, size); + for (int i = 0; i < height; i++) { + for (int j = 0; j < width; j++) { + hData[i*width+j] = i*width+j; + } + } + printf("hData: "); + for (int i = 0; i < 64; i++) { + printf("%f ", hData[i]); + } + printf("\n"); + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat); + hipArray *hipArray; + hipMallocArray(&hipArray, &channelDesc, width, height); + + hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice); + + struct hipResourceDesc resDesc; + memset(&resDesc, 0, sizeof(resDesc)); + resDesc.resType = hipResourceTypeArray; + resDesc.res.array.array = hipArray; + + // Specify texture object parameters + struct hipTextureDesc texDesc; + memset(&texDesc, 0, sizeof(texDesc)); + texDesc.addressMode[0] = hipAddressModeWrap; + texDesc.addressMode[1] = hipAddressModeWrap; + texDesc.filterMode = hipFilterModePoint; + texDesc.readMode = hipReadModeElementType; + texDesc.normalizedCoords = 0; + + // Create texture object + hipTextureObject_t textureObject = 0; + hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL); + + float* dData = NULL; + hipMalloc((void **) &dData, size); + + dim3 dimBlock(16, 16, 1); + dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1); + + hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, textureObject, width, height); + + hipDeviceSynchronize(); + + float *hOutputData = (float *) malloc(size); + memset(hOutputData, 0, size); + hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost); + + printf("dData: "); + for (int i = 0; i < 64; i++) { + printf("%f ", hOutputData[i]); + } + printf("\n"); + for (int i = 0; i < height; i++) { + for (int j = 0; j < width; j++) { + if (hData[i*width+j] != hOutputData[i*width+j]) { + printf("Difference [ %d %d ]:%f ----%f\n",i, j, hData[i*width+j] , hOutputData[i*width+j]); + testResult = false; + break; + } + } + } + hipDestroyTextureObject(textureObject); + hipFree(dData); + hipFreeArray(hipArray); +} diff --git a/tests/src/texture/hipTextureRef2D.cpp b/tests/src/texture/hipTextureRef2D.cpp new file mode 100644 index 0000000000..04537a0a83 --- /dev/null +++ b/tests/src/texture/hipTextureRef2D.cpp @@ -0,0 +1,100 @@ +#include +#include +#include + +#include +#include "test_common.h" +texture tex; + +bool testResult = true; + +__global__ void tex2DKernel(float* outputData, + hipTextureObject_t textureObject, + int width, + int height) +{ + int x = hipBlockIdx_x*hipBlockDim_x + hipThreadIdx_x; + int y = hipBlockIdx_y*hipBlockDim_y + hipThreadIdx_y; +#ifdef __HIP_PLATFORM_HCC__ + outputData[y*width + x] = tex2D(tex, textureObject, x, y); +#else + outputData[y*width + x] = tex2D(tex, x, y); +#endif +} + +void runTest(int argc, char **argv); + +int main(int argc, char **argv) +{ + runTest(argc, argv); + if(testResult) { + passed(); + } else { + exit(EXIT_FAILURE); + } +} + +void runTest(int argc, char **argv) +{ + unsigned int width = 256; + unsigned int height = 256; + unsigned int size = width * height * sizeof(float); + float* hData = (float*) malloc(size); + memset(hData, 0, size); + for (int i = 0; i < height; i++) { + for (int j = 0; j < width; j++) { + hData[i*width+j] = i*width+j; + } + } + printf("hData: "); + for (int i = 0; i < 64; i++) { + printf("%f ", hData[i]); + } + printf("\n"); + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat); + hipArray *hipArray; + hipMallocArray(&hipArray, &channelDesc, width, height); + + hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice); + + tex.addressMode[0] = hipAddressModeWrap; + tex.addressMode[1] = hipAddressModeWrap; + tex.filterMode = hipFilterModePoint; + tex.normalized = 0; + + hipBindTextureToArray(tex, hipArray, channelDesc); + + float* dData = NULL; + hipMalloc((void **) &dData, size); + + dim3 dimBlock(16, 16, 1); + dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1); +#ifdef __HIP_PLATFORM_HCC__ + hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, tex.textureObject, width, height); +#else + hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, 0, width, height); +#endif + hipDeviceSynchronize(); + + float *hOutputData = (float *) malloc(size); + memset(hOutputData, 0, size); + hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost); + + printf("dData: "); + for (int i = 0; i < 64; i++) { + printf("%f ", hOutputData[i]); + } + printf("\n"); + for (int i = 0; i < height; i++) { + for (int j = 0; j < width; j++) { + if (hData[i*width+j] != hOutputData[i*width+j]) { + printf("Difference [ %d %d ]:%f ----%f\n",i, j, hData[i*width+j] , hOutputData[i*width+j]); + testResult = false; + break; + } + } + } + hipFree(dData); + hipFreeArray(hipArray); +} From ce13e4afbfd4c1221aa53d87ace0adbc7d722a61 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 26 Sep 2017 11:39:09 +0530 Subject: [PATCH 2/8] Added HIT block --- tests/src/texture/hipTextureObj2D.cpp | 5 +++++ tests/src/texture/hipTextureRef2D.cpp | 5 +++++ 2 files changed, 10 insertions(+) diff --git a/tests/src/texture/hipTextureObj2D.cpp b/tests/src/texture/hipTextureObj2D.cpp index bea21d0658..443d708418 100644 --- a/tests/src/texture/hipTextureObj2D.cpp +++ b/tests/src/texture/hipTextureObj2D.cpp @@ -1,3 +1,8 @@ +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * RUN: %t + * HIT_END + */ #include #include #include diff --git a/tests/src/texture/hipTextureRef2D.cpp b/tests/src/texture/hipTextureRef2D.cpp index 04537a0a83..eb27b23230 100644 --- a/tests/src/texture/hipTextureRef2D.cpp +++ b/tests/src/texture/hipTextureRef2D.cpp @@ -1,3 +1,8 @@ +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * RUN: %t + * HIT_END + */ #include #include #include From 470224cf2585e5d7bf2d59b31ba0d8f384b9b985 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Sat, 23 Sep 2017 05:23:24 +0530 Subject: [PATCH 3/8] Fix texture support on HIP/NVCC path --- include/hip/nvcc_detail/hip_runtime_api.h | 174 ++++++++++++++++++++-- 1 file changed, 162 insertions(+), 12 deletions(-) diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h index 1103bb2274..6459839180 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/include/hip/nvcc_detail/hip_runtime_api.h @@ -50,15 +50,64 @@ hipMemcpyHostToHost ,hipMemcpyDeviceToDevice ,hipMemcpyDefault } hipMemcpyKind ; +#if 0 +typedef enum hipTextureAddressMode { + hipAddressModeWrap = 0, + hipAddressModeClamp = 1, + hipAddressModeMirror = 2, + hipAddressModeBorder = 3 +}hipTextureAddressMode; +#endif +#define hipTextureAddressMode cudaTextureAddressMode +#define hipAddressModeWrap cudaAddressModeWrap +#define hipAddressModeClamp cudaAddressModeClamp +#define hipAddressModeMirror cudaAddressModeMirror +#define hipAddressModeBorder cudaAddressModeBorder +#if 0 +typedef enum hipTextureFilterMode { + hipFilterModePoint = 0, + hipFilterModeLinear = 1 +}hipTextureFilterMode; + +#endif +#if 1 +#define hipTextureFilterMode cudaTextureFilterMode +#define hipFilterModePoint cudaFilterModePoint +#define hipFilterModeLinear cudaFilterModeLinear +#endif +#if 0 +typedef enum hipTextureReadMode { + hipReadModeElementType = 0, + hipReadModeNormalizedFloat = 1 +}hipTextureReadMode; +#endif +#define hipTextureReadMode cudaTextureReadMode +#define hipReadModeElementType cudaReadModeElementType +#define hipReadModeNormalizedFloat cudaReadModeNormalizedFloat + +typedef enum hipChannelFormatKind { + hipChannelFormatKindSigned = 0, + hipChannelFormatKindUnsigned = 1, + hipChannelFormatKindFloat = 2, + hipChannelFormatKindNone = 3 +}hipChannelFormatKind; +#if 0 +typedef enum hipResourceType { + hipResourceTypeArray = 0, + hipResourceTypeMipmappedArray = 1, + hipResourceTypeLinear = 2, + hipResourceTypePitch2D = 3 +}hipResourceType; +#endif +#define hipResourceType cudaResourceType +#define hipResourceTypeArray cudaResourceTypeArray +#define hipResourceTypeMipmappedArray cudaResourceTypeMipmappedArray +#define hipResourceTypeLinear cudaResourceTypeLinear +#define hipResourceTypePitch2D cudaResourceTypePitch2D +// // hipErrorNoDevice. -/*typedef enum hipTextureFilterMode -{ - hipFilterModePoint = cudaFilterModePoint, ///< Point filter mode. -//! @warning cudaFilterModeLinear is not supported. -} hipTextureFilterMode;*/ -#define hipFilterModePoint cudaFilterModePoint //! Flags that can be used with hipEventCreateWithFlags: #define hipEventDefault cudaEventDefault @@ -119,10 +168,16 @@ typedef CUdevice hipDevice_t; typedef CUmodule hipModule_t; typedef CUfunction hipFunction_t; typedef CUdeviceptr hipDeviceptr_t; -typedef enum cudaChannelFormatKind hipChannelFormatKind; -typedef struct cudaChannelFormatDesc hipChannelFormatDesc; -typedef enum cudaTextureReadMode hipTextureReadMode; +//typedef enum cudaChannelFormatKind hipChannelFormatKind; +//typedef struct cudaChannelFormatDesc hipChannelFormatDesc; typedef struct cudaArray hipArray; +typedef struct cudaArray* hipArray_const_t; +//typedef cudaArray_const_t hipArray; +#define hipArrayDefault cudaArrayDefault + +typedef cudaTextureObject_t hipTextureObject_t; +#define hipTextureType2D cudaTextureType2D; +#define hipDeviceMapHost cudaDeviceMapHost // Flags that can be used with hipStreamCreateWithFlags #define hipStreamDefault cudaStreamDefault @@ -130,7 +185,9 @@ typedef struct cudaArray hipArray; //typedef cudaChannelFormatDesc hipChannelFormatDesc; #define hipChannelFormatDesc cudaChannelFormatDesc - +#define hipResourceDesc cudaResourceDesc +#define hipTextureDesc cudaTextureDesc +#define hipResourceViewDesc cudaResourceViewDesc //adding code for hipmemSharedConfig #define hipSharedMemBankSizeDefault cudaSharedMemBankSizeDefault #define hipSharedMemBankSizeFourByte cudaSharedMemBankSizeFourByte @@ -214,6 +271,58 @@ inline static enum cudaMemcpyKind hipMemcpyKindToCudaMemcpyKind(hipMemcpyKind ki } } +inline static cudaTextureAddressMode hipTextureAddressModeToCudaTextureAddressMode(hipTextureAddressMode kind) { + switch(kind) { + case hipAddressModeWrap: + return cudaAddressModeWrap; + case hipAddressModeClamp: + return cudaAddressModeClamp; + case hipAddressModeMirror: + return cudaAddressModeMirror; + case hipAddressModeBorder: + return cudaAddressModeBorder; + default: + return cudaAddressModeWrap; +} +} + +inline static cudaTextureFilterMode hipTextureFilterModeToCudaTextureFilterMode(hipTextureFilterMode kind) { + switch(kind) { + case hipFilterModePoint: + return cudaFilterModePoint; + case hipFilterModeLinear: + return cudaFilterModeLinear; + default: + return cudaFilterModePoint; +} +} + +inline static cudaTextureReadMode hipTextureReadModeToCudaTextureReadMode(hipTextureReadMode kind) { + switch(kind) { + case hipReadModeElementType: + return cudaReadModeElementType; + case hipReadModeNormalizedFloat: + return cudaReadModeNormalizedFloat; + default: + return cudaReadModeElementType; +} +} + +inline static cudaChannelFormatKind hipChannelFormatKindToCudaChannelFormatKind(hipChannelFormatKind kind) { + switch(kind) { + case hipChannelFormatKindSigned : + return cudaChannelFormatKindSigned ; + case hipChannelFormatKindUnsigned : + return cudaChannelFormatKindUnsigned ; + case hipChannelFormatKindFloat : + return cudaChannelFormatKindFloat ; + case hipChannelFormatKindNone : + return cudaChannelFormatKindNone ; + default: + return cudaChannelFormatKindNone ; +} +} + /** * Stream CallBack struct */ @@ -262,9 +371,17 @@ inline static hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int fla return hipCUDAErrorTohipError(cudaHostAlloc(ptr, size, flags)); } -inline static hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc, size_t width, size_t height, unsigned int flags) { - return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags)); +#if __cplusplus +inline static hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, + size_t width, size_t height = 0, unsigned int flags = hipArrayDefault) { + return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags)); } +#else +inline static hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc, + size_t width, size_t height, unsigned int flags) { + return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags)); +} +#endif inline static hipError_t hipFreeArray(hipArray* array) { return hipCUDAErrorTohipError(cudaFreeArray(array)); @@ -1007,11 +1124,44 @@ inline static hipError_t hipUnbindTexture(struct texture *tex return hipCUDAErrorTohipError(cudaUnbindTexture(tex)); } +template +inline static hipError_t hipBindTextureToArray(struct texture& tex, + hipArray_const_t array, + const struct hipChannelFormatDesc& desc) +{ + return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array,desc)); +} + +template +inline static hipError_t hipBindTextureToArray(struct texture &tex, + hipArray_const_t array) +{ + return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array)); +} + template inline static hipChannelFormatDesc hipCreateChannelDesc() { return cudaCreateChannelDesc(); } + +hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f) +{ + return cudaCreateChannelDesc(x , y , z , w, hipChannelFormatKindToCudaChannelFormatKind(f)); +} + +inline static hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, + const hipResourceDesc* pResDesc, + const hipTextureDesc* pTexDesc, + const hipResourceViewDesc* pResViewDesc) +{ + return hipCUDAErrorTohipError(cudaCreateTextureObject(pTexObject, pResDesc, pTexDesc, pResViewDesc)); +} + +inline static hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject) +{ + return hipCUDAErrorTohipError(cudaDestroyTextureObject(textureObject)); +} #endif //__CUDACC__ #endif //HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_RUNTIME_API_H From f62254a8d85c9ed74a546130433514371117fb34 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Mon, 25 Sep 2017 17:32:56 +0530 Subject: [PATCH 4/8] Fixed build issue --- include/hip/nvcc_detail/hip_runtime_api.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h index 6459839180..d8d0bc2c62 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/include/hip/nvcc_detail/hip_runtime_api.h @@ -1145,7 +1145,7 @@ inline static hipChannelFormatDesc hipCreateChannelDesc() return cudaCreateChannelDesc(); } -hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f) +inline static hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f) { return cudaCreateChannelDesc(x , y , z , w, hipChannelFormatKindToCudaChannelFormatKind(f)); } From 19d574d939419e830e2c56aecf7d5f8d9234540f Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Mon, 2 Oct 2017 10:57:25 +0530 Subject: [PATCH 5/8] Fixed review comments --- include/hip/nvcc_detail/hip_runtime_api.h | 46 +++++------------------ 1 file changed, 9 insertions(+), 37 deletions(-) diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h index d8d0bc2c62..245327a041 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/include/hip/nvcc_detail/hip_runtime_api.h @@ -50,38 +50,20 @@ hipMemcpyHostToHost ,hipMemcpyDeviceToDevice ,hipMemcpyDefault } hipMemcpyKind ; -#if 0 -typedef enum hipTextureAddressMode { - hipAddressModeWrap = 0, - hipAddressModeClamp = 1, - hipAddressModeMirror = 2, - hipAddressModeBorder = 3 -}hipTextureAddressMode; -#endif + +//hipTextureAddressMode #define hipTextureAddressMode cudaTextureAddressMode #define hipAddressModeWrap cudaAddressModeWrap #define hipAddressModeClamp cudaAddressModeClamp #define hipAddressModeMirror cudaAddressModeMirror #define hipAddressModeBorder cudaAddressModeBorder -#if 0 -typedef enum hipTextureFilterMode { - hipFilterModePoint = 0, - hipFilterModeLinear = 1 -}hipTextureFilterMode; - -#endif -#if 1 +//hipTextureFilterMode #define hipTextureFilterMode cudaTextureFilterMode #define hipFilterModePoint cudaFilterModePoint #define hipFilterModeLinear cudaFilterModeLinear -#endif -#if 0 -typedef enum hipTextureReadMode { - hipReadModeElementType = 0, - hipReadModeNormalizedFloat = 1 -}hipTextureReadMode; -#endif + +//hipTextureReadMode #define hipTextureReadMode cudaTextureReadMode #define hipReadModeElementType cudaReadModeElementType #define hipReadModeNormalizedFloat cudaReadModeNormalizedFloat @@ -92,14 +74,8 @@ typedef enum hipChannelFormatKind { hipChannelFormatKindFloat = 2, hipChannelFormatKindNone = 3 }hipChannelFormatKind; -#if 0 -typedef enum hipResourceType { - hipResourceTypeArray = 0, - hipResourceTypeMipmappedArray = 1, - hipResourceTypeLinear = 2, - hipResourceTypePitch2D = 3 -}hipResourceType; -#endif + +//hipResourceType #define hipResourceType cudaResourceType #define hipResourceTypeArray cudaResourceTypeArray #define hipResourceTypeMipmappedArray cudaResourceTypeMipmappedArray @@ -168,11 +144,8 @@ typedef CUdevice hipDevice_t; typedef CUmodule hipModule_t; typedef CUfunction hipFunction_t; typedef CUdeviceptr hipDeviceptr_t; -//typedef enum cudaChannelFormatKind hipChannelFormatKind; -//typedef struct cudaChannelFormatDesc hipChannelFormatDesc; typedef struct cudaArray hipArray; typedef struct cudaArray* hipArray_const_t; -//typedef cudaArray_const_t hipArray; #define hipArrayDefault cudaArrayDefault typedef cudaTextureObject_t hipTextureObject_t; @@ -183,7 +156,6 @@ typedef cudaTextureObject_t hipTextureObject_t; #define hipStreamDefault cudaStreamDefault #define hipStreamNonBlocking cudaStreamNonBlocking -//typedef cudaChannelFormatDesc hipChannelFormatDesc; #define hipChannelFormatDesc cudaChannelFormatDesc #define hipResourceDesc cudaResourceDesc #define hipTextureDesc cudaTextureDesc @@ -371,14 +343,14 @@ inline static hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int fla return hipCUDAErrorTohipError(cudaHostAlloc(ptr, size, flags)); } -#if __cplusplus +#if 0 //__cplusplus inline static hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, size_t width, size_t height = 0, unsigned int flags = hipArrayDefault) { return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags)); } #else inline static hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc, - size_t width, size_t height, unsigned int flags) { + size_t width, size_t height, unsigned int flags __dparm(hipArrayDefault)) { return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags)); } #endif From b8aae66424ba170a3a16d04373f58d7105954541 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Mon, 2 Oct 2017 17:29:08 +0530 Subject: [PATCH 6/8] Fixed review comment --- include/hip/nvcc_detail/hip_runtime_api.h | 7 ------- 1 file changed, 7 deletions(-) diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h index 245327a041..870dcc5dae 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/include/hip/nvcc_detail/hip_runtime_api.h @@ -343,17 +343,10 @@ inline static hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int fla return hipCUDAErrorTohipError(cudaHostAlloc(ptr, size, flags)); } -#if 0 //__cplusplus -inline static hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, - size_t width, size_t height = 0, unsigned int flags = hipArrayDefault) { - return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags)); -} -#else inline static hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc, size_t width, size_t height, unsigned int flags __dparm(hipArrayDefault)) { return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags)); } -#endif inline static hipError_t hipFreeArray(hipArray* array) { return hipCUDAErrorTohipError(cudaFreeArray(array)); From 68c2146dc3a498051e8111217094da407612bd9a Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 2 Oct 2017 21:52:16 +0000 Subject: [PATCH 7/8] Fix math ordering for --genco mode. --- include/hip/hcc_detail/math_functions.h | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/include/hip/hcc_detail/math_functions.h b/include/hip/hcc_detail/math_functions.h index 79bacf274b..efc15ea70c 100644 --- a/include/hip/hcc_detail/math_functions.h +++ b/include/hip/hcc_detail/math_functions.h @@ -23,6 +23,11 @@ THE SOFTWARE. #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_MATH_FUNCTIONS_H #define HIP_INCLUDE_HIP_HCC_DETAIL_MATH_FUNCTIONS_H +#if defined(__HCC__) + #include +#endif + + #include #include #include From ca239ab6ef79060914eef5403a1e8e9550acd941 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 6 Oct 2017 10:26:11 +0530 Subject: [PATCH 8/8] Bump HIP version to 1.4 and update release notes Change-Id: I3570dedb32fe5fec6f60565c587a3282a4c6c709 --- RELEASE.md | 11 +++++++++++ bin/hipconfig | 2 +- 2 files changed, 12 insertions(+), 1 deletion(-) diff --git a/RELEASE.md b/RELEASE.md index 452ac54ee1..3987255f04 100644 --- a/RELEASE.md +++ b/RELEASE.md @@ -8,6 +8,17 @@ We have attempted to document known bugs and limitations - in particular the [HI ## Revision History: +=================================================================================================== +Release: 1.4 +Date: 2017.10.06 +- Improvements to HIP event management +- Added new HIP_TRACE_API options +- Enabled device side assert support +- Several bug fixes including hipMallocArray, hipTexture fetch +- Support for RHEL/CentOS 7.4 +- Updates to hipify-perl, hipify-clang and documentation + + =================================================================================================== Release: 1.3 Date: 2017.08.16 diff --git a/bin/hipconfig b/bin/hipconfig index c74d757fb5..cd8d36c5e7 100755 --- a/bin/hipconfig +++ b/bin/hipconfig @@ -1,7 +1,7 @@ #!/usr/bin/perl -w $HIP_BASE_VERSION_MAJOR = "1"; -$HIP_BASE_VERSION_MINOR = "3"; +$HIP_BASE_VERSION_MINOR = "4"; # Need perl > 5.10 to use logic-defined or use 5.006; use v5.10.1;