diff --git a/hipamd/include/hip/channel_descriptor.h b/hipamd/include/hip/channel_descriptor.h new file mode 100644 index 0000000000..af8875e256 --- /dev/null +++ b/hipamd/include/hip/channel_descriptor.h @@ -0,0 +1,36 @@ +/* +Copyright (c) 2015-2017 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 + +// 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: + + +#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__) +#include +#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__) +#include +#else +#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__"); +#endif diff --git a/hipamd/include/hip/hcc_detail/channel_descriptor.h b/hipamd/include/hip/hcc_detail/channel_descriptor.h new file mode 100644 index 0000000000..91d48af4cf --- /dev/null +++ b/hipamd/include/hip/hcc_detail/channel_descriptor.h @@ -0,0 +1,374 @@ +/* +Copyright (c) 2015-2017 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. +*/ + +#ifndef HIP_HCC_DETAIL_CHANNEL_DESCRIPTOR_H +#define HIP_HCC_DETAIL_CHANNEL_DESCRIPTOR_H + +#include +#include + +hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f); + +template +static inline hipChannelFormatDesc hipCreateChannelDesc() { + return hipCreateChannelDesc(0, 0, 0, 0, hipChannelFormatKindNone); +} + +static inline hipChannelFormatDesc hipCreateChannelDescHalf() { + int e = (int)sizeof(unsigned short) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindFloat); +} + +static inline hipChannelFormatDesc hipCreateChannelDescHalf1() { + int e = (int)sizeof(unsigned short) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindFloat); +} + +static inline hipChannelFormatDesc hipCreateChannelDescHalf2() +{ + int e = (int)sizeof(unsigned short) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindFloat); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(char) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(signed char) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(unsigned char) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(unsigned char) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(signed char) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(unsigned char) * 8; + return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(signed char) * 8; + return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(unsigned char) * 8; + return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(signed char) * 8; + return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(unsigned char) * 8; + return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(signed char) * 8; + return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(unsigned short) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(signed short) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(unsigned short) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(signed short) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(unsigned short) * 8; + return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindUnsigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(signed short) * 8; + return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(unsigned short) * 8; + return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindUnsigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(signed short) * 8; + return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(unsigned short) * 8; + return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindUnsigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(signed short) * 8; + return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(unsigned int) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(signed int) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(unsigned int) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(signed int) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(unsigned int) * 8; + return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindUnsigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(signed int) * 8; + return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(unsigned int) * 8; + return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindUnsigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(signed int) * 8; + return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(unsigned int) * 8; + return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindUnsigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(signed int) * 8; + return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(float) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindFloat); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(float) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindFloat); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(float) * 8; + return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindFloat); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(float) * 8; + return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindFloat); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(float) * 8; + return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindFloat); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(unsigned long) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(signed long) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(unsigned long) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(signed long) * 8; + return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(unsigned long) * 8; + return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindUnsigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(signed long) * 8; + return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(unsigned long) * 8; + return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindUnsigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(signed long) * 8; + return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindSigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(unsigned long) * 8; + return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindUnsigned); +} + +template<> +inline hipChannelFormatDesc hipCreateChannelDesc() +{ + int e = (int)sizeof(signed long) * 8; + return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindSigned); +} + +#endif diff --git a/hipamd/include/hip/hcc_detail/driver_types.h b/hipamd/include/hip/hcc_detail/driver_types.h new file mode 100644 index 0000000000..1fe72b0507 --- /dev/null +++ b/hipamd/include/hip/hcc_detail/driver_types.h @@ -0,0 +1,41 @@ +/* +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. +*/ + +#ifndef HIP_HCC_DETAIL_DRIVER_TYPES_H +#define HIP_HCC_DETAIL_DRIVER_TYPES_H + +enum hipChannelFormatKind +{ + hipChannelFormatKindSigned = 0, + hipChannelFormatKindUnsigned = 1, + hipChannelFormatKindFloat = 2, + hipChannelFormatKindNone = 3 +}; + +struct hipChannelFormatDesc +{ + int x; + int y; + int z; + int w; + enum hipChannelFormatKind f; +}; + + +#endif diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 448102a1c6..f9ab03901a 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015-2017 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 @@ -21,8 +21,8 @@ THE SOFTWARE. */ //#pragma once -#ifndef HIP_RUNTIME_API_H -#define HIP_RUNTIME_API_H +#ifndef HIP_HCC_DETAIL_HIP_RUNTIME_API_H +#define HIP_HCC_DETAIL_HIP_RUNTIME_API_H /** * @file hcc_detail/hip_runtime_api.h * @brief Contains C function APIs for HIP runtime. This file does not use any HCC builtin or special language extensions (-hc mode) ; those functions in hip_runtime.h. @@ -177,6 +177,12 @@ typedef enum hipMemcpyKind { ,hipMemcpyDefault = 4, ///< Runtime will automatically determine copy-kind based on virtual addresses. } hipMemcpyKind; +typedef struct { + unsigned int width; + unsigned int height; + hipChannelFormatKind f; + void* data; //FIXME: generalize this +} hipArray; @@ -1235,6 +1241,84 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t st **/ hipError_t hipMemGetInfo (size_t * free, size_t * total) ; + +/** + * @brief Allocate an array on the device. + * + * @param[out] array Pointer to allocated array in device memory + * @param[in] desc Requested channel format + * @param[in] width Requested array allocation width + * @param[in] height Requested array allocation height + * @param[in] flags Requested properties of allocated array + * @return #hipSuccess, #hipErrorMemoryAllocation + * + * @see hipMalloc, hipMallocPitch, hipFree, hipFreeArray, hipHostMalloc, hipHostFree + */ +hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, + size_t width, size_t height = 0, unsigned int flags = 0); + +/** + * @brief Frees an array on the device. + * + * @param[in] array Pointer to array to free + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInitializationError + * + * @see hipMalloc, hipMallocPitch, hipFree, hipMallocArray, hipHostMalloc, hipHostFree + */ +hipError_t hipFreeArray(hipArray* array); + +/** + * @brief Copies data between host and device. + * + * @param[in] dst Destination memory address + * @param[in] dpitch Pitch of destination memory + * @param[in] src Source memory address + * @param[in] spitch Pitch of source memory + * @param[in] width Width of matrix transfer (columns in bytes) + * @param[in] height Height of matrix transfer (rows) + * @param[in] kind Type of transfer + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection + * + * @see hipMemcpy, hipMemcpyToArray, hipMemcpy2DToArray, hipMemcpyFromArray, hipMemcpyToSymbol, hipMemcpyAsync + */ +hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind); + +/** + * @brief Copies data between host and device. + * + * @param[in] dst Destination memory address + * @param[in] dpitch Pitch of destination memory + * @param[in] src Source memory address + * @param[in] spitch Pitch of source memory + * @param[in] width Width of matrix transfer (columns in bytes) + * @param[in] height Height of matrix transfer (rows) + * @param[in] kind Type of transfer + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection + * + * @see hipMemcpy, hipMemcpyToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol, hipMemcpyAsync + */ +hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, + size_t spitch, size_t width, size_t height, hipMemcpyKind kind); + +/** + * @brief Copies data between host and device. + * + * @param[in] dst Destination memory address + * @param[in] dpitch Pitch of destination memory + * @param[in] src Source memory address + * @param[in] spitch Pitch of source memory + * @param[in] width Width of matrix transfer (columns in bytes) + * @param[in] height Height of matrix transfer (rows) + * @param[in] kind Type of transfer + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection + * + * @see hipMemcpy, hipMemcpy2DToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol, hipMemcpyAsync + */ +hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, + const void* src, size_t count, hipMemcpyKind kind); + + + // doxygen end Memory /** * @} @@ -1911,6 +1995,87 @@ hipError_t hipIpcCloseMemHandle(void *devPtr); } /* extern "c" */ #endif +#ifdef __cplusplus +/* + * @brief hipBindTexture Binds size bytes of the memory area pointed to by @p devPtr to the texture reference tex. + * + * @p desc describes how the memory is interpreted when fetching values from the texture. The @p offset parameter is an optional byte offset as with the low-level + * hipBindTexture() function. Any memory previously bound to tex is unbound. + * + * @param[in] offset - Offset in bytes + * @param[out] tex - texture to bind + * @param[in] devPtr - Memory area on device + * @param[in] desc - Channel format + * @param[in] size - Size of the memory area pointed to by devPtr + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree, #hipErrorUnknown + **/ +template +hipError_t hipBindTexture(size_t *offset, + struct texture &tex, + const void *devPtr, + const struct hipChannelFormatDesc *desc, + size_t size=UINT_MAX) +{ + tex._dataPtr = static_cast(devPtr); + + return hipSuccess; +} + +/* + * @brief hipBindTexture Binds size bytes of the memory area pointed to by @p devPtr to the texture reference tex. + * + * @p desc describes how the memory is interpreted when fetching values from the texture. The @p offset parameter is an optional byte offset as with the low-level + * hipBindTexture() function. Any memory previously bound to tex is unbound. + * + * @param[in] offset - Offset in bytes + * @param[in] tex - texture to bind + * @param[in] devPtr - Memory area on device + * @param[in] size - Size of the memory area pointed to by devPtr + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree, #hipErrorUnknown + **/ +template +hipError_t hipBindTexture(size_t *offset, + struct texture &tex, + const void *devPtr, + size_t size=UINT_MAX) +{ + return hipBindTexture(offset, tex, devPtr, &tex.channelDesc, size); +} + +template +hipError_t hipBindTextureToArray(struct texture &tex, hipArray* array) { + tex.width = array->width; + tex.height = array->height; + tex._dataPtr = static_cast(array->data); + return hipSuccess; +} + +/* + * @brief Unbinds the textuer bound to @p tex + * + * @param[in] tex - texture to unbind + * + * @return #hipSuccess + **/ +template +hipError_t hipUnbindTexture(struct texture &tex) +{ + tex._dataPtr = NULL; + + return hipSuccess; +} + + + +// doxygen end Texture +/** + * @} + */ + + +#endif + + /** *------------------------------------------------------------------------------------------------- *------------------------------------------------------------------------------------------------- diff --git a/hipamd/include/hip/hcc_detail/hip_texture.h b/hipamd/include/hip/hcc_detail/hip_texture.h index fe13f11e49..bd44206be4 100644 --- a/hipamd/include/hip/hcc_detail/hip_texture.h +++ b/hipamd/include/hip/hcc_detail/hip_texture.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015-2017 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 @@ -31,48 +31,16 @@ THE SOFTWARE. */ #include - +#include +#include //#include //---- //Texture - TODO - likely need to move this to a separate file only included with kernel compilation. #define hipTextureType1D 1 -typedef enum { - hipChannelFormatKindSigned = 0, - hipChannelFormatKindUnsigned, - hipChannelFormatKindFloat, - hipChannelFormatKindNone - -} hipChannelFormatKind; - -typedef struct hipChannelFormatDesc { - int x; - int y; - int z; - int w; - hipChannelFormatKind f; -} 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; -}; #if __cplusplus -template +template struct texture : public textureReference { const T * _dataPtr; // pointer to underlying data. @@ -84,95 +52,12 @@ struct texture : public textureReference { }; #endif -typedef struct { - unsigned int width; - unsigned int height; - hipChannelFormatKind f; - void* data; //FIXME: generalize this -} hipArray; - #define tex1Dfetch(_tex, _addr) (_tex._dataPtr[_addr]) #define tex2D(_tex, _dx, _dy) \ _tex._dataPtr[(unsigned int)_dx + (unsigned int)_dy*(_tex.width)] -/** - * @brief Allocate an array on the device. - * - * @param[out] array Pointer to allocated array in device memory - * @param[in] desc Requested channel format - * @param[in] width Requested array allocation width - * @param[in] height Requested array allocation height - * @param[in] flags Requested properties of allocated array - * @return #hipSuccess, #hipErrorMemoryAllocation - * - * @see hipMalloc, hipMallocPitch, hipFree, hipFreeArray, hipHostMalloc, hipHostFree - */ -hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, - size_t width, size_t height = 0, unsigned int flags = 0); - -/** - * @brief Frees an array on the device. - * - * @param[in] array Pointer to array to free - * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInitializationError - * - * @see hipMalloc, hipMallocPitch, hipFree, hipMallocArray, hipHostMalloc, hipHostFree - */ -hipError_t hipFreeArray(hipArray* array); - -/** - * @brief Copies data between host and device. - * - * @param[in] dst Destination memory address - * @param[in] dpitch Pitch of destination memory - * @param[in] src Source memory address - * @param[in] spitch Pitch of source memory - * @param[in] width Width of matrix transfer (columns in bytes) - * @param[in] height Height of matrix transfer (rows) - * @param[in] kind Type of transfer - * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection - * - * @see hipMemcpy, hipMemcpyToArray, hipMemcpy2DToArray, hipMemcpyFromArray, hipMemcpyToSymbol, hipMemcpyAsync - */ -hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind); - -/** - * @brief Copies data between host and device. - * - * @param[in] dst Destination memory address - * @param[in] dpitch Pitch of destination memory - * @param[in] src Source memory address - * @param[in] spitch Pitch of source memory - * @param[in] width Width of matrix transfer (columns in bytes) - * @param[in] height Height of matrix transfer (rows) - * @param[in] kind Type of transfer - * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection - * - * @see hipMemcpy, hipMemcpyToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol, hipMemcpyAsync - */ -hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, - size_t spitch, size_t width, size_t height, hipMemcpyKind kind); - -/** - * @brief Copies data between host and device. - * - * @param[in] dst Destination memory address - * @param[in] dpitch Pitch of destination memory - * @param[in] src Source memory address - * @param[in] spitch Pitch of source memory - * @param[in] width Width of matrix transfer (columns in bytes) - * @param[in] height Height of matrix transfer (rows) - * @param[in] kind Type of transfer - * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection - * - * @see hipMemcpy, hipMemcpy2DToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol, hipMemcpyAsync - */ -hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, - const void* src, size_t count, hipMemcpyKind kind); - - /** * @addtogroup API HIP API * @{ @@ -212,120 +97,6 @@ hipChannelFormatDesc hipBindTexture(size_t *offset, struct textureReference *te } #endif -/** - * @brief Returns a channel descriptor using the specified format. - * - * @param[in] x X component - * @param[in] y Y component - * @param[in] z Z component - * @param[in] w W component - * @param[in] f Channel format - * @return Channel descriptor with format f - * - */ -hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f); - -// descriptors -template inline hipChannelFormatDesc hipCreateChannelDesc() { - return hipCreateChannelDesc(0, 0, 0, 0, hipChannelFormatKindNone); -} -template <> inline hipChannelFormatDesc hipCreateChannelDesc() { - int e = (int)sizeof(int) * 8; - return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned); -} -template <> inline hipChannelFormatDesc hipCreateChannelDesc() { - int e = (int)sizeof(unsigned int) * 8; - return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned); -} -template <> inline hipChannelFormatDesc hipCreateChannelDesc() { - int e = (int)sizeof(long) * 8; - return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned); -} -template <> inline hipChannelFormatDesc hipCreateChannelDesc() { - int e = (int)sizeof(unsigned long) * 8; - return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned); -} -template <> inline hipChannelFormatDesc hipCreateChannelDesc() { - int e = (int)sizeof(float) * 8; - return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindFloat); -} - -/* - * @brief hipBindTexture Binds size bytes of the memory area pointed to by @p devPtr to the texture reference tex. - * - * @p desc describes how the memory is interpreted when fetching values from the texture. The @p offset parameter is an optional byte offset as with the low-level - * hipBindTexture() function. Any memory previously bound to tex is unbound. - * - * @param[in] offset - Offset in bytes - * @param[out] tex - texture to bind - * @param[in] devPtr - Memory area on device - * @param[in] desc - Channel format - * @param[in] size - Size of the memory area pointed to by devPtr - * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree, #hipErrorUnknown - **/ -template -hipError_t hipBindTexture(size_t *offset, - struct texture &tex, - const void *devPtr, - const struct hipChannelFormatDesc *desc, - size_t size=UINT_MAX) -{ - tex._dataPtr = static_cast(devPtr); - - return hipSuccess; -} - -/* - * @brief hipBindTexture Binds size bytes of the memory area pointed to by @p devPtr to the texture reference tex. - * - * @p desc describes how the memory is interpreted when fetching values from the texture. The @p offset parameter is an optional byte offset as with the low-level - * hipBindTexture() function. Any memory previously bound to tex is unbound. - * - * @param[in] offset - Offset in bytes - * @param[in] tex - texture to bind - * @param[in] devPtr - Memory area on device - * @param[in] size - Size of the memory area pointed to by devPtr - * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree, #hipErrorUnknown - **/ -template -hipError_t hipBindTexture(size_t *offset, - struct texture &tex, - const void *devPtr, - size_t size=UINT_MAX) -{ - return hipBindTexture(offset, tex, devPtr, &tex.channelDesc, size); -} - -template -hipError_t hipBindTextureToArray(struct texture &tex, hipArray* array) { - tex.width = array->width; - tex.height = array->height; - tex._dataPtr = static_cast(array->data); - return hipSuccess; -} - -/* - * @brief Unbinds the textuer bound to @p tex - * - * @param[in] tex - texture to unbind - * - * @return #hipSuccess - **/ -template -hipError_t hipUnbindTexture(struct texture &tex) -{ - tex._dataPtr = NULL; - - return hipSuccess; -} - - - -// doxygen end Texture -/** - * @} - */ - // End doxygen API: /** @@ -333,4 +104,3 @@ hipError_t hipUnbindTexture(struct texture &tex) */ #endif - diff --git a/hipamd/include/hip/hcc_detail/texture_types.h b/hipamd/include/hip/hcc_detail/texture_types.h new file mode 100644 index 0000000000..d9221dc009 --- /dev/null +++ b/hipamd/include/hip/hcc_detail/texture_types.h @@ -0,0 +1,42 @@ +/* +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. +*/ + + +#ifndef HIP_HCC_DETAIL_TEXTURE_TYPES_H +#define HIP_HCC_DETAIL_TEXTURE_TYPES_H + +#include + +enum hipTextureReadMode +{ + hipReadModeElementType = 0 +}; + +enum hipTextureFilterMode +{ + hipFilterModePoint = 0 +}; + +struct textureReference { + hipTextureFilterMode filterMode; + bool normalized; + hipChannelFormatDesc channelDesc; +}; + +#endif diff --git a/hipamd/include/hip/math_functions.h b/hipamd/include/hip/math_functions.h index d33f7a2e90..6afe463a5d 100644 --- a/hipamd/include/hip/math_functions.h +++ b/hipamd/include/hip/math_functions.h @@ -20,19 +20,6 @@ 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 diff --git a/hipamd/include/hip/nvcc_detail/channel_descriptor.h b/hipamd/include/hip/nvcc_detail/channel_descriptor.h new file mode 100644 index 0000000000..8502745968 --- /dev/null +++ b/hipamd/include/hip/nvcc_detail/channel_descriptor.h @@ -0,0 +1,25 @@ +/* +Copyright (c) 2015-2017 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"channel_descriptor.h" diff --git a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 482e91e440..3cd5667175 100644 --- a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015-2017 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