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; 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 diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h index 1103bb2274..870dcc5dae 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/include/hip/nvcc_detail/hip_runtime_api.h @@ -51,14 +51,39 @@ hipMemcpyHostToHost ,hipMemcpyDefault } hipMemcpyKind ; +//hipTextureAddressMode +#define hipTextureAddressMode cudaTextureAddressMode +#define hipAddressModeWrap cudaAddressModeWrap +#define hipAddressModeClamp cudaAddressModeClamp +#define hipAddressModeMirror cudaAddressModeMirror +#define hipAddressModeBorder cudaAddressModeBorder + +//hipTextureFilterMode +#define hipTextureFilterMode cudaTextureFilterMode +#define hipFilterModePoint cudaFilterModePoint +#define hipFilterModeLinear cudaFilterModeLinear + +//hipTextureReadMode +#define hipTextureReadMode cudaTextureReadMode +#define hipReadModeElementType cudaReadModeElementType +#define hipReadModeNormalizedFloat cudaReadModeNormalizedFloat + +typedef enum hipChannelFormatKind { + hipChannelFormatKindSigned = 0, + hipChannelFormatKindUnsigned = 1, + hipChannelFormatKindFloat = 2, + hipChannelFormatKindNone = 3 +}hipChannelFormatKind; + +//hipResourceType +#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,18 +144,22 @@ 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 struct cudaArray hipArray; +typedef struct cudaArray* hipArray_const_t; +#define hipArrayDefault cudaArrayDefault + +typedef cudaTextureObject_t hipTextureObject_t; +#define hipTextureType2D cudaTextureType2D; +#define hipDeviceMapHost cudaDeviceMapHost // Flags that can be used with hipStreamCreateWithFlags #define hipStreamDefault cudaStreamDefault #define hipStreamNonBlocking cudaStreamNonBlocking -//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 +243,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,8 +343,9 @@ 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)); +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)); } inline static hipError_t hipFreeArray(hipArray* array) { @@ -1007,11 +1089,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(); } + +inline static 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 diff --git a/tests/src/texture/hipTextureObj2D.cpp b/tests/src/texture/hipTextureObj2D.cpp new file mode 100644 index 0000000000..443d708418 --- /dev/null +++ b/tests/src/texture/hipTextureObj2D.cpp @@ -0,0 +1,112 @@ +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * RUN: %t + * HIT_END + */ +#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..eb27b23230 --- /dev/null +++ b/tests/src/texture/hipTextureRef2D.cpp @@ -0,0 +1,105 @@ +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * RUN: %t + * HIT_END + */ +#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); +}