From bce37a03b8178676603ba6c94ad86119a15bf4bb Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Thu, 16 May 2019 20:36:51 +0000 Subject: [PATCH 01/10] Fix hipcc exit code when failing [ROCm/hip commit: 20ea7e9d26b2a2db8ea8d31ca14b426da928b1a8] --- projects/hip/bin/hipcc | 16 +++++++++++++++- 1 file changed, 15 insertions(+), 1 deletion(-) diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index fadc74c137..012545e5d4 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -848,7 +848,21 @@ if ($runCmd) { print ("HIP ($HIP_PATH) was built using hcc $hipConfig{'HCC_VERSION'}, but you are using $HCC_HOME/hcc with version $HCC_VERSION from hipcc. Please rebuild HIP including cmake or update HCC_HOME variable.\n") ; die unless $ENV{'HIP_IGNORE_HCC_VERSION'}; } - system ("$CMD") or delete_temp_dirs () and die (); + system ("$CMD"); + if ($? == -1) { + print "failed to execute: $!\n"; + exit($?); + } + elsif ($? & 127) { + printf "child died with signal %d, %s coredump\n", + ($? & 127), ($? & 128) ? 'with' : 'without'; + exit($?); + } + else { + $CMD_EXIT_CODE = $? >> 8; + } + $? or delete_temp_dirs (); + exit($CMD_EXIT_CODE); } # vim: ts=4:sw=4:expandtab:smartindent From 993f323887d0c7fa007d0731b21a77bc34ac6770 Mon Sep 17 00:00:00 2001 From: Konstantin Pyzhov Date: Mon, 20 May 2019 02:13:30 -0400 Subject: [PATCH 02/10] Disable hipStreamSync2 test for hip-clang. [ROCm/hip commit: 7cbaa5082b978b87cc361e1cd0b894c332cbc12c] --- projects/hip/tests/src/runtimeApi/stream/hipStreamSync2.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamSync2.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamSync2.cpp index c365dce70a..8160ac4bc3 100644 --- a/projects/hip/tests/src/runtimeApi/stream/hipStreamSync2.cpp +++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamSync2.cpp @@ -178,7 +178,9 @@ void runTests(int64_t numElements) { { test(0x01, C_d, C_h, numElements, syncNone, true /*expectMismatch*/); test(0x02, C_d, C_h, numElements, syncNullStream, false /*expectMismatch*/); +#ifndef __HIP_CLANG_ONLY__ test(0x04, C_d, C_h, numElements, syncOtherStream, true /*expectMismatch*/); +#endif test(0x08, C_d, C_h, numElements, syncDevice, false /*expectMismatch*/); // Sending a marker to to null stream may synchronize the otherStream From a4202d5ce6b4c213426177c1c16ea006279fdab5 Mon Sep 17 00:00:00 2001 From: Konstantin Pyzhov Date: Tue, 21 May 2019 13:21:48 -0400 Subject: [PATCH 03/10] Updated kernel binary file for hipModule test. [ROCm/hip commit: a8b07ada2f730461f388d79da8f13191950192e4] --- .../src/runtimeApi/module/vcpy_kernel.code | Bin 9456 -> 9456 bytes 1 file changed, 0 insertions(+), 0 deletions(-) diff --git a/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.code b/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.code index 737ee1ce9b3eedb45e5cf76377c851dccf2aca11..4246151be15f18f6ca4c6924779f4d6b05680982 100755 GIT binary patch delta 186 zcmez1`N4C-46b=>3}7&iZR4UkM$QIi2)ALfCQ~!hU#7|HnBHvWWASEWGBleU&H8eZ z0Lx|{wk@2K4cIFt33xCHO!gGm&B(BsRgjsVX$R|ME(JNpjLDh`@74nb4CD+ CI4BDM delta 210 zcmez1`N4C-3@$Nt1~3p~-?*rbk@E{Pg!^T(CQ~z0A@k&QOm7%XCeLP;0+InNQtVEZ zVTn1Vsgv(8OEG&I8*Of2PGw{=GMW6G<>kZ)ESoQ|?% Date: Wed, 22 May 2019 18:31:39 +0300 Subject: [PATCH 04/10] [HIP][HIPIFY] Make hipMemcpyParam2D coherent with cuMemcpy2D + Makes hip_Memcpy2D struct compatible with CUDA_MEMCPY2D struct + Add hipMemcpyParam2D support in nvcc fallback path + Update hipify-clang, tests and docs accordingly [ROCm/hip commit: 8f059b0ee928d2d2f25c03c81d9f0dac3b1d5f62] --- ...A_Driver_API_functions_supported_by_HIP.md | 2 +- .../src/CUDA2HIP_Driver_API_functions.cpp | 4 +-- .../hip/include/hip/hcc_detail/driver_types.h | 28 +++++++++---------- .../include/hip/hcc_detail/hip_runtime_api.h | 10 +++++++ .../include/hip/nvcc_detail/hip_runtime_api.h | 12 ++++++++ .../11_texture_driver/texture2dDrv.cpp | 4 +-- projects/hip/src/hip_memory.cpp | 4 +-- 7 files changed, 43 insertions(+), 21 deletions(-) diff --git a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index aeb834c2ec..7c2d2c6631 100644 --- a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md @@ -861,7 +861,7 @@ | `cuMemAllocManaged` | | | `cuMemAllocPitch` | | | `cuMemcpy` | | -| `cuMemcpy2D` | | +| `cuMemcpy2D` | `hipMemcpyParam2D` | | `cuMemcpy2DAsync` | | | `cuMemcpy2DUnaligned` | | | `cuMemcpy3D` | | diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp index b71e49710d..8be20774ea 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp @@ -175,8 +175,8 @@ const std::map CUDA_DRIVER_FUNCTION_MAP{ {"cuMemcpy", {"hipMemcpy_", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, // no analogue // NOTE: Not equal to cudaMemcpy2D due to different signatures - {"cuMemcpy2D", {"hipMemcpy2D_", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, - {"cuMemcpy2D_v2", {"hipMemcpy2D_", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuMemcpy2D", {"hipMemcpyParam2D", "", CONV_MEMORY, API_DRIVER}}, + {"cuMemcpy2D_v2", {"hipMemcpyParam2D", "", CONV_MEMORY, API_DRIVER}}, // no analogue // NOTE: Not equal to cudaMemcpy2DAsync due to different signatures {"cuMemcpy2DAsync", {"hipMemcpy2DAsync_", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, diff --git a/projects/hip/include/hip/hcc_detail/driver_types.h b/projects/hip/include/hip/hcc_detail/driver_types.h index 8e1fec11fa..5b2297114f 100644 --- a/projects/hip/include/hip/hcc_detail/driver_types.h +++ b/projects/hip/include/hip/hcc_detail/driver_types.h @@ -80,22 +80,22 @@ typedef struct hipArray { }hipArray; 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; + hipMemoryType srcMemoryType; + const void* srcHost; + hipDeviceptr_t srcDevice; + hipArray* srcArray; + size_t srcPitch; + size_t dstXInBytes; + size_t dstY; + hipMemoryType dstMemoryType; + void* dstHost; + hipDeviceptr_t dstDevice; + hipArray* dstArray; + size_t dstPitch; + size_t WidthInBytes; + size_t Height; } hip_Memcpy2D; diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index 5b598b54a8..d870963101 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -1858,6 +1858,16 @@ hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc* */ hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind); + +/** +* @brief Copies memory for 2D arrays. +* @param[in] pCopy Parameters for the memory copy + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, + * #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection + * + * @see hipMemcpy, hipMemcpy2D, hipMemcpyToArray, hipMemcpy2DToArray, hipMemcpyFromArray, + * hipMemcpyToSymbol, hipMemcpyAsync +*/ hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy); /** diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index 2ebd4f8b7d..08897fdb07 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -145,6 +145,12 @@ typedef enum hipChannelFormatKind { #define hipJitOptionFastCompile CU_JIT_FAST_COMPILE #define hipJitOptionNumOptions CU_JIT_NUM_OPTIONS +// enum CUmemorytype redefines +#define hipMemoryTypeHost CU_MEMORYTYPE_HOST +#define hipMemoryTypeDevice CU_MEMORYTYPE_DEVICE +#define hipMemoryTypeArray CU_MEMORYTYPE_ARRAY +#define hipMemoryTypeUnified CU_MEMORYTYPE_UNIFIED + typedef cudaEvent_t hipEvent_t; typedef cudaStream_t hipStream_t; typedef cudaIpcEventHandle_t hipIpcEventHandle_t; @@ -162,6 +168,8 @@ typedef CUdeviceptr hipDeviceptr_t; typedef struct cudaArray hipArray; typedef struct cudaArray* hipArray_const_t; typedef cudaFuncAttributes hipFuncAttributes; +typedef enum CUmemorytype hipMemoryType; +typedef struct CUDA_MEMCPY2D hip_Memcpy2D; #define hipMemcpy3DParms cudaMemcpy3DParms #define hipArrayDefault cudaArrayDefault #define hipArrayLayered cudaArrayLayered @@ -578,6 +586,10 @@ inline static hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, cudaMemcpy2D(dst, dpitch, src, spitch, width, height, hipMemcpyKindToCudaMemcpyKind(kind))); } +inline static hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) { + return hipCUDAErrorTohipError(cuMemcpy2D(hip_Memcpy2D)); +} + inline static hipError_t hipMemcpy3D(const struct hipMemcpy3DParms *p) { return hipCUDAErrorTohipError(cudaMemcpy3D(p)); diff --git a/projects/hip/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp b/projects/hip/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp index a6ba44696a..36d37a4fad 100644 --- a/projects/hip/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp +++ b/projects/hip/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp @@ -71,8 +71,8 @@ bool runTest(int argc, char** argv) { copyParam.srcMemoryType = hipMemoryTypeHost; copyParam.srcHost = hData; copyParam.srcPitch = width * sizeof(float); - copyParam.widthInBytes = copyParam.srcPitch; - copyParam.height = height; + copyParam.WidthInBytes = copyParam.srcPitch; + copyParam.Height = height; hipMemcpyParam2D(©Param); textureReference* texref; diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 0fad8ab890..36edcdb338 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -1715,8 +1715,8 @@ hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) { if (pCopy == nullptr) { e = hipErrorInvalidValue; } - e = ihipMemcpy2D(pCopy->dstArray->data, pCopy->widthInBytes, pCopy->srcHost, pCopy->srcPitch, - pCopy->widthInBytes, pCopy->height, hipMemcpyDefault); + e = ihipMemcpy2D(pCopy->dstArray->data, pCopy->WidthInBytes, pCopy->srcHost, pCopy->srcPitch, + pCopy->WidthInBytes, pCopy->Height, hipMemcpyDefault); return ihipLogStatus(e); } From 9ebc77525456c9d8cae6fdcfc4bda40295c478b1 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 22 May 2019 20:04:45 +0300 Subject: [PATCH 05/10] [HIP] fix nvcc path break in #1127 [ROCm/hip commit: 47bbb9c605bce9a84c9bcedbee740649a6a85442] --- .../hip/include/hip/nvcc_detail/hip_runtime_api.h | 11 ++--------- 1 file changed, 2 insertions(+), 9 deletions(-) diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index 08897fdb07..7bc7b91368 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -145,12 +145,6 @@ typedef enum hipChannelFormatKind { #define hipJitOptionFastCompile CU_JIT_FAST_COMPILE #define hipJitOptionNumOptions CU_JIT_NUM_OPTIONS -// enum CUmemorytype redefines -#define hipMemoryTypeHost CU_MEMORYTYPE_HOST -#define hipMemoryTypeDevice CU_MEMORYTYPE_DEVICE -#define hipMemoryTypeArray CU_MEMORYTYPE_ARRAY -#define hipMemoryTypeUnified CU_MEMORYTYPE_UNIFIED - typedef cudaEvent_t hipEvent_t; typedef cudaStream_t hipStream_t; typedef cudaIpcEventHandle_t hipIpcEventHandle_t; @@ -168,8 +162,7 @@ typedef CUdeviceptr hipDeviceptr_t; typedef struct cudaArray hipArray; typedef struct cudaArray* hipArray_const_t; typedef cudaFuncAttributes hipFuncAttributes; -typedef enum CUmemorytype hipMemoryType; -typedef struct CUDA_MEMCPY2D hip_Memcpy2D; +#define hip_Memcpy2D CUDA_MEMCPY2D #define hipMemcpy3DParms cudaMemcpy3DParms #define hipArrayDefault cudaArrayDefault #define hipArrayLayered cudaArrayLayered @@ -587,7 +580,7 @@ inline static hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, } inline static hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) { - return hipCUDAErrorTohipError(cuMemcpy2D(hip_Memcpy2D)); + return hipCUResultTohipError(cuMemcpy2D(hip_Memcpy2D)); } inline static hipError_t hipMemcpy3D(const struct hipMemcpy3DParms *p) From c24d90afb3298dadf439340822808ded45a4207f Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 22 May 2019 20:48:18 +0300 Subject: [PATCH 06/10] [HIP] fix typo in #1127 [ROCm/hip commit: eede980a398e37aebcd7e6a715b532a1d39b3163] --- projects/hip/include/hip/nvcc_detail/hip_runtime_api.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index 7bc7b91368..c1846c1b1e 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -580,7 +580,7 @@ inline static hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, } inline static hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) { - return hipCUResultTohipError(cuMemcpy2D(hip_Memcpy2D)); + return hipCUResultTohipError(cuMemcpy2D(pCopy)); } inline static hipError_t hipMemcpy3D(const struct hipMemcpy3DParms *p) From b08cbff718d998213a6204e7aa772d51b6490f3e Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Wed, 22 May 2019 19:20:09 +0000 Subject: [PATCH 07/10] Fix bug in __smid not setting correct size The SZ field should minus by 1 since SIZE range is 1..32. Also add comments that results may vary. [ROCm/hip commit: ff74babedfd74b671166d3db3420bc091b1938e8] --- projects/hip/include/hip/hcc_detail/device_functions.h | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/device_functions.h b/projects/hip/include/hip/hcc_detail/device_functions.h index e1f81e4547..808ed216e4 100644 --- a/projects/hip/include/hip/hcc_detail/device_functions.h +++ b/projects/hip/include/hip/hcc_detail/device_functions.h @@ -1009,14 +1009,20 @@ void __syncthreads() #define GETREG_IMMED(SZ,OFF,REG) (SZ << 11) | (OFF << 6) | REG +/* + __smid returns the wave's assigned Compute Unit and Shader Engine. + The Compute Unit, CU_ID returned in bits 3:0, and Shader Engine, SE_ID in bits 5:4. + Note: the results vary over time. + SZ minus 1 since SIZE is 1-based. +*/ __device__ inline unsigned __smid(void) { unsigned cu_id = __builtin_amdgcn_s_getreg( - GETREG_IMMED(HW_ID_CU_ID_SIZE, HW_ID_CU_ID_OFFSET, HW_ID)); + GETREG_IMMED(HW_ID_CU_ID_SIZE-1, HW_ID_CU_ID_OFFSET, HW_ID)); unsigned se_id = __builtin_amdgcn_s_getreg( - GETREG_IMMED(HW_ID_SE_ID_SIZE, HW_ID_SE_ID_OFFSET, HW_ID)); + GETREG_IMMED(HW_ID_SE_ID_SIZE-1, HW_ID_SE_ID_OFFSET, HW_ID)); /* Each shader engine has 16 CU */ return (se_id << HW_ID_CU_ID_SIZE) + cu_id; From b0d091506c2588124dc20b9bd41a777c77cd0709 Mon Sep 17 00:00:00 2001 From: Brian Sumner Date: Wed, 22 May 2019 12:19:51 -0700 Subject: [PATCH 08/10] Update kernel language documentation [ROCm/hip commit: 2d782664d118b65f03e848574cbc9c858d7d3e1a] --- .../hip/docs/markdown/hip_kernel_language.md | 67 +++++++++++++++++++ 1 file changed, 67 insertions(+) diff --git a/projects/hip/docs/markdown/hip_kernel_language.md b/projects/hip/docs/markdown/hip_kernel_language.md index d69f5a04a8..5479813675 100644 --- a/projects/hip/docs/markdown/hip_kernel_language.md +++ b/projects/hip/docs/markdown/hip_kernel_language.md @@ -35,6 +35,9 @@ - [Warp Cross-Lane Functions](#warp-cross-lane-functions) * [Warp Vote and Ballot Functions](#warp-vote-and-ballot-functions) * [Warp Shuffle Functions](#warp-shuffle-functions) +- [Cooperative Groups Functions](#cooperative-groups-functions) +- [Warp Matrix Functions](#warp-matrix-functions) +- [Independent Thread Scheduling](#independent-thread-scheduling) - [Profiler Counter Function](#profiler-counter-function) - [Assert](#assert) - [Printf](#printf) @@ -599,6 +602,70 @@ float __shfl_xor (float var, int laneMask, int width=warpSize); ``` +## Cooperative Groups Functions + +Cooperative groups is a mechanism for forming and communicating between groups of threads at +a granularity different than the block. This feature was introduced in Cuda 9. + +HIP does not support any of the kernel language cooperative groups +types or functions. + + +| **Function** | **Supported in HIP** | **Supported in CUDA** | +| --- | --- | --- | +| `void thread_group.sync()` | | ✓ | +| `unsigned thread_group.size()` | | ✓ | +| `unsigned thread_group.thread_rank()` | | ✓ | +| `bool thread_group.is_valid()` | | ✓ | +| `thread_group tiled_partition(thread_group, size)` | | ✓ | +| `thread_block_tile tiled_partition(thread_group)` | | ✓ | +| `thread_block this_thread_block()` | | ✓ | +| `T thread_block_tile.shfl()` | | ✓ | +| `T thread_block_tile.shfl_down()` | | ✓ | +| `T thread_block_tile.shfl_up()` | | ✓ | +| `T thread_block_tile.shfl_xor()` | | ✓ | +| `T thread_block_tile.any()` | | ✓ | +| `T thread_block_tile.all()` | | ✓ | +| `T thread_block_tile.ballot()` | | ✓ | +| `T thread_block_tile.match_any()` | | ✓ | +| `T thread_block_tile.match_all()` | | ✓ | +| `coalesced_group coalesced_threads()` | | ✓ | +| `grid_group this_grid()` | | ✓ | +| `void grid_group.sync()` | | ✓ | +| `unsigned grid_group.size()` | | ✓ | +| `unsigned grid_group.thread_rank()` | | ✓ | +| `bool grid_group.is_valid()` | | ✓ | +| `multi_grid_group this_multi_grid()` | | ✓ | +| `void multi_grid_group.sync()` | | ✓ | +| `unsigned multi_grid_group.size()` | | ✓ | +| `unsigned multi_grid_group.thread_rank()` | | ✓ | +| `bool multi_grid_group.is_valid()` | | ✓ | + +## Warp Matrix Functions + +Warp matrix functions allow a warp to cooperatively operate on small matrices +whose elements are spread over the lanes in an unspecified manner. This feature +was introduced in Cuda 9. + +HIP does not support any of the kernel language warp matrix +types or functions. + +| **Function** | **Supported in HIP** | **Supported in CUDA** | +| --- | --- | --- | +| `void load_matrix_sync(fragment<...> &a, const T* mptr, unsigned lda)` | | ✓ | +| `void load_matrix_sync(fragment<...> &a, const T* mptr, unsigned lda, layout_t layout)` | | ✓ | +| `void store_matrix_sync(T* mptr, fragment<...> &a, unsigned lda, layout_t layout)` | | ✓ | +| `void fill_fragment(fragment<...> &a, const T &value)` | | ✓ | +| `void mma_sync(fragment<...> &d, const fragment<...> &a, const fragment<...> &b, const fragment<...> &c , bool sat)` | | ✓ | + +## Independent Thread Scheduling + +The hardware support for independent thread scheduling introduced in certain architectures +supporting Cuda allows threads to progress independently of each other and enables +intra-warp synchronizations that were previously not allowed. + +HIP does not support this type of scheduling. + ## Profiler Counter Function The Cuda `__prof_trigger()` instruction is not supported. From 6ba67c4a8110d81b0a76ae12cc7519a23a0a2248 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 23 May 2019 12:47:08 +0300 Subject: [PATCH 09/10] [HIPIFY] Add device texture fetch functions support + Add a corresponding reverse engineered sample tex2dKernel with texture template [ROCm/hip commit: c7a2b4e492785e38e5ae7ba5df282728275feaeb] --- projects/hip/hipify-clang/src/CUDA2HIP.cpp | 1 + .../hip/hipify-clang/src/HipifyAction.cpp | 3 ++ .../11_texture_driver/tex2dKernel.cpp | 36 +++++++++++++++++++ 3 files changed, 40 insertions(+) create mode 100644 projects/hip/tests/hipify-clang/unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp diff --git a/projects/hip/hipify-clang/src/CUDA2HIP.cpp b/projects/hip/hipify-clang/src/CUDA2HIP.cpp index 9a26a051ca..1e530745e6 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP.cpp @@ -33,6 +33,7 @@ const std::map CUDA_INCLUDE_MAP{ {"driver_types.h", {"hip/driver_types.h", "", CONV_INCLUDE, API_RUNTIME}}, {"cuda_fp16.h", {"hip/hip_fp16.h", "", CONV_INCLUDE, API_RUNTIME}}, {"cuda_texture_types.h", {"hip/hip_texture_types.h", "", CONV_INCLUDE, API_RUNTIME}}, + {"texture_fetch_functions.h", {"", "", CONV_INCLUDE, API_RUNTIME}}, {"vector_types.h", {"hip/hip_vector_types.h", "", CONV_INCLUDE, API_RUNTIME}}, {"cuda_profiler_api.h", {"hip/hip_profile.h", "", CONV_INCLUDE, API_RUNTIME}}, // cuComplex includes diff --git a/projects/hip/hipify-clang/src/HipifyAction.cpp b/projects/hip/hipify-clang/src/HipifyAction.cpp index 5ea83c8376..241ca7ecae 100644 --- a/projects/hip/hipify-clang/src/HipifyAction.cpp +++ b/projects/hip/hipify-clang/src/HipifyAction.cpp @@ -189,6 +189,9 @@ bool HipifyAction::Exclude(const hipCounter & hipToken) { } return false; case CONV_INCLUDE: + if (hipToken.hipName.empty()) { + return true; + } switch (hipToken.apiType) { case API_RAND: if (hipToken.hipName == "hiprand_kernel.h") { diff --git a/projects/hip/tests/hipify-clang/unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp b/projects/hip/tests/hipify-clang/unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp new file mode 100644 index 0000000000..d5dffd0b09 --- /dev/null +++ b/projects/hip/tests/hipify-clang/unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp @@ -0,0 +1,36 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +/* +Copyright (c) 2015-present 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. +*/ + +// CHECK: #include +#include +// CHECK-NOT: #include +#include + +// CHECK: extern texture tex; +extern texture tex; + +extern "C" __global__ void tex2dKernel(float* outputData, int width, int height) { + int x = blockDim.x * blockIdx.x + threadIdx.x; + int y = blockDim.y * blockIdx.y + threadIdx.y; + outputData[y * width + x] = tex2D(tex, x, y); +} From 320676a2e3eeb49ff1e3b62918eb4a5f985cac41 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Thu, 23 May 2019 15:11:02 -0400 Subject: [PATCH 10/10] Fix device lib path for hip-clang. We now always need device lib path when compiling and not need it at linking. [ROCm/hip commit: bba85f0e24d7baf7d0b45500b611c674efda01b7] --- projects/hip/bin/hipcc | 15 +-------------- 1 file changed, 1 insertion(+), 14 deletions(-) diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index fadc74c137..e46fe5604f 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -370,7 +370,6 @@ if($HIP_PLATFORM eq "nvcc"){ my $toolArgs = ""; # arguments to pass to the hcc or nvcc tool my $optArg = ""; # -O args -my $rdc = 0; foreach $arg (@ARGV) { @@ -466,14 +465,6 @@ foreach $arg (@ARGV) { $optArg = $arg; } - if($arg =~ /-fgpu-rdc/) - { - $rdc = 1; - } - if($arg =~ /-fno-gpu-rdc/) - { - $rdc = 0; - } ## process linker response file for hip-clang ## extract object files from static library and pass them directly to @@ -804,11 +795,7 @@ if ($HIP_PLATFORM eq "clang") { $HIPLDFLAGS .= " -O3"; } $HIP_DEVLIB_FLAGS = " --hip-device-lib-path=$DEVICE_LIB_PATH"; - if ($rdc eq 0) { - $HIPCXXFLAGS .= " $HIP_DEVLIB_FLAGS"; - } else { - $HIPLDFLAGS .= " $HIP_DEVLIB_FLAGS"; - } + $HIPCXXFLAGS .= " $HIP_DEVLIB_FLAGS"; if ($isWindows) { $HIPCXXFLAGS .= " -std=c++14 -fms-extensions -fms-compatibility"; } else {