Texture driver APIs support

[ROCm/hip commit: ef09c4918d]
2017-11-09 22:10:55 +05:30
parent 56d3677499
commit 49a2c6468e
8 changed files with 4114 additions and 17 deletions
@@ -23,6 +23,7 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_DRIVER_TYPES_H
 #define HIP_INCLUDE_HIP_HCC_DETAIL_DRIVER_TYPES_H

+typedef void* hipDeviceptr_t;
 enum hipChannelFormatKind
 {
    hipChannelFormatKindSigned = 0,
@@ -40,6 +41,29 @@ struct hipChannelFormatDesc
    enum hipChannelFormatKind f;
 };

+#define HIP_TRSF_NORMALIZED_COORDINATES 0x02
+#define HIP_TRSF_READ_AS_INTEGER 0x01
+#define HIP_TRSA_OVERRIDE_FORMAT 0x01
+
+enum hipArray_Format
+{
+    HIP_AD_FORMAT_UNSIGNED_INT8 = 0x01,
+    HIP_AD_FORMAT_UNSIGNED_INT16 = 0x02,
+    HIP_AD_FORMAT_UNSIGNED_INT32 = 0x03,
+    HIP_AD_FORMAT_SIGNED_INT8 = 0x08,
+    HIP_AD_FORMAT_SIGNED_INT16 = 0x09,
+    HIP_AD_FORMAT_SIGNED_INT32 = 0x0a,
+    HIP_AD_FORMAT_HALF = 0x10,
+    HIP_AD_FORMAT_FLOAT = 0x20
+};
+
+struct HIP_ARRAY_DESCRIPTOR {
+	enum hipArray_Format format;
+	unsigned int numChannels;
+	size_t width;
+	size_t height;
+};
+
 struct hipArray {
    void* data; //FIXME: generalize this
    struct hipChannelFormatDesc desc;
@@ -47,8 +71,30 @@ struct hipArray {
    unsigned int width;
    unsigned int height;
    unsigned int depth;
+    struct HIP_ARRAY_DESCRIPTOR drvDesc;
+    bool isDrv;
 };

+typedef struct hip_Memcpy2D {
+    size_t  height;
+    size_t  widthInBytes;
+    hipArray* dstArray;
+    hipDeviceptr_t dstDevice;
+    void * dstHost;
+    hipMemoryType dstMemoryType;
+    size_t  dstPitch;
+    size_t  dstXInBytes;
+    size_t  dstY;
+    hipArray* srcArray;
+    hipDeviceptr_t srcDevice;
+    const void * srcHost;
+    hipMemoryType srcMemoryType;
+    size_t  srcPitch;
+    size_t  srcXInBytes;
+    size_t  srcY;
+}hip_Memcpy2D;
+
+
 typedef struct hipArray* hipArray_t;

 typedef const struct hipArray* hipArray_const_t;
@@ -84,8 +84,6 @@ typedef struct ihipModule_t *hipModule_t;

 typedef struct ihipModuleSymbol_t *hipFunction_t;

-typedef void* hipDeviceptr_t;
-
 typedef struct ihipEvent_t *hipEvent_t;

 enum hipLimit_t
@@ -621,7 +619,7 @@ hipError_t hipStreamQuery(hipStream_t stream);
 *
 * This command is host-synchronous : the host will block until the specified stream is empty.
 *
- * This command follows standard null-stream semantics.  Specifically, specifying the null stream will cause the 
+ * This command follows standard null-stream semantics.  Specifically, specifying the null stream will cause the
 * command to wait for other streams on the same device to complete all pending operations.
 *
 * This command honors the hipDeviceLaunchBlocking flag, which controls whether the wait is active or blocking.
@@ -644,9 +642,9 @@ hipError_t hipStreamSynchronize(hipStream_t stream);
 * This function inserts a wait operation into the specified stream.
 * All future work submitted to @p stream will wait until @p event reports completion before beginning execution.
 *
- * This function only waits for commands in the current stream to complete.  Notably,, this function does 
- * not impliciy wait for commands in the default stream to complete, even if the specified stream is 
- * created with hipStreamNonBlocking = 0.  
+ * This function only waits for commands in the current stream to complete.  Notably,, this function does
+ * not impliciy wait for commands in the default stream to complete, even if the specified stream is
+ * created with hipStreamNonBlocking = 0.
 *
 * @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamSynchronize, hipStreamDestroy
 */
@@ -756,7 +754,7 @@ hipError_t hipEventCreate(hipEvent_t* event);
 * If hipEventRecord() has been previously called on this event, then this call will overwrite any existing state in event.
 *
 * If this function is called on a an event that is currently being recorded, results are undefined - either
- * outstanding recording may save state into the event, and the order is not guaranteed.  
+ * outstanding recording may save state into the event, and the order is not guaranteed.
 *
 * @see hipEventCreate, hipEventCreateWithFlags, hipEventQuery, hipEventSynchronize, hipEventDestroy, hipEventElapsedTime
 *
@@ -1318,6 +1316,7 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc,
 hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc,
                          size_t width, size_t height, unsigned int flags);
 #endif
+hipError_t hipArrayCreate ( hipArray** pHandle, const HIP_ARRAY_DESCRIPTOR* pAllocateArray );
 /**
 *  @brief Frees an array on the device.
 *
@@ -1359,6 +1358,7 @@ hipError_t hipMalloc3DArray(hipArray_t *array,
 *  @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);
+hipError_t hipMemcpy_2D(const hip_Memcpy2D* pCopy);

 /**
 *  @brief Copies data between host and device.
@@ -1968,6 +1968,7 @@ hipError_t hipModuleGetFunction(hipFunction_t *function, hipModule_t module, con
 */
 hipError_t hipModuleGetGlobal(hipDeviceptr_t *dptr, size_t *bytes, hipModule_t hmod, const char *name);

+hipError_t hipModuleGetTexRef(hipDeviceptr_t *dptr, hipModule_t hmod, const char* name);
 /**
 * @brief builds module from code object which resides in host memory. Image is pointer to that location.
 *
@@ -2359,6 +2360,15 @@ hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject);
 hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDesc, hipTextureObject_t textureObject);
 hipError_t hipGetTextureObjectResourceViewDesc(hipResourceViewDesc* pResViewDesc, hipTextureObject_t textureObject);
 hipError_t hipGetTextureObjectTextureDesc(hipTextureDesc* pTexDesc, hipTextureObject_t textureObject);
+hipError_t hipTexRefSetArray ( textureReference* tex,  hipArray_const_t array, unsigned int  flags );
+
+hipError_t hipTexRefSetAddressMode ( textureReference* tex, int  dim, hipTextureAddressMode am );
+
+hipError_t hipTexRefSetFilterMode ( textureReference*  tex, hipTextureFilterMode fm );
+
+hipError_t hipTexRefSetFlags ( textureReference*  tex, unsigned int  flags );
+
+hipError_t hipTexRefSetFormat (textureReference* tex, hipArray_Format fmt, int  NumPackedComponents );

 // doxygen end Texture
 /**
@@ -93,6 +93,8 @@ struct textureReference
    float                       maxMipmapLevelClamp;

    hipTextureObject_t          textureObject;
+    int 						numChannels;
+    enum hipArray_Format        format;
 };

 /**
@@ -117,7 +117,9 @@ typedef struct hipDeviceProp_t {
 */
 enum hipMemoryType {
    hipMemoryTypeHost,   ///< Memory is physically located on host
-    hipMemoryTypeDevice  ///< Memory is physically located on device. (see deviceId for specific device)
+    hipMemoryTypeDevice,  ///< Memory is physically located on device. (see deviceId for specific device)
+    hipMemoryTypeArray,
+    hipMemoryTypeUnified
 };