From 8ba092c58f1a349f2915c6b7aad2044c994568c1 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 30 Jun 2022 11:38:36 +0530 Subject: [PATCH 1/3] SWDEV-321698 - hipGraphNodeTypeMemcpy1D NVDIA build failure (#2531) * SWDEV-321698 - hipGraphNodeTypeMemcpy1D NVDIA build failure Signed-off-by: sdashmiz Change-Id: Ibb3c7278438cd6504f9bd4035fe7cc27d0cd6c0f * Update hip_runtime_api.h * Update hip_runtime_api.h [ROCm/hip commit: 64a4f67f5fd899e8551091ffb09a58bce133c1f4] --- projects/hip/include/hip/hip_runtime_api.h | 23 +++++++++++----------- 1 file changed, 12 insertions(+), 11 deletions(-) diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index c6706024bc..83e2638fdd 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -1082,17 +1082,18 @@ typedef struct hipGraphExec* hipGraphExec_t; * */ typedef enum hipGraphNodeType { - hipGraphNodeTypeKernel = 1, ///< GPU kernel node - hipGraphNodeTypeMemcpy = 2, ///< Memcpy 3D node - hipGraphNodeTypeMemset = 3, ///< Memset 1D node - hipGraphNodeTypeHost = 4, ///< Host (executable) node - hipGraphNodeTypeGraph = 5, ///< Node which executes an embedded graph - hipGraphNodeTypeEmpty = 6, ///< Empty (no-op) node - hipGraphNodeTypeWaitEvent = 7, ///< External event wait node - hipGraphNodeTypeEventRecord = 8, ///< External event record node - hipGraphNodeTypeMemcpy1D = 9, ///< Memcpy 1D node - hipGraphNodeTypeMemcpyFromSymbol = 10, ///< MemcpyFromSymbol node - hipGraphNodeTypeMemcpyToSymbol = 11, ///< MemcpyToSymbol node + hipGraphNodeTypeKernel = 0, ///< GPU kernel node + hipGraphNodeTypeMemcpy = 1, ///< Memcpy node + hipGraphNodeTypeMemset = 2, ///< Memset node + hipGraphNodeTypeHost = 3, ///< Host (executable) node + hipGraphNodeTypeGraph = 4, ///< Node which executes an embedded graph + hipGraphNodeTypeEmpty = 5, ///< Empty (no-op) node + hipGraphNodeTypeWaitEvent = 6, ///< External event wait node + hipGraphNodeTypeEventRecord = 7, ///< External event record node + hipGraphNodeTypeExtSemaphoreSignal = 8, ///< External Semaphore signal node + hipGraphNodeTypeExtSemaphoreWait = 9, ///< External Semaphore wait node + hipGraphNodeTypeMemcpyFromSymbol = 10, ///< MemcpyFromSymbol node + hipGraphNodeTypeMemcpyToSymbol = 11, ///< MemcpyToSymbol node hipGraphNodeTypeCount } hipGraphNodeType; From 47a1198e182405cf241d7f65900f2c20f25ca26b Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 30 Jun 2022 13:28:15 +0530 Subject: [PATCH 2/3] SWDEV-342662 - Added hipGraphUpload API (#2757) Change-Id: I65a957e4d8b12b2ffd003e5a4722edaed2ffa99e [ROCm/hip commit: 84c23e2b41a384ec771c489e39a4f34e096d7ea1] --- projects/hip/include/hip/hip_runtime_api.h | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index 83e2638fdd..f8a244d046 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -5618,6 +5618,18 @@ hipError_t hipGraphInstantiateWithFlags(hipGraphExec_t* pGraphExec, hipGraph_t g */ hipError_t hipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream); +/** + * @brief uploads an executable graph in a stream + * + * @param [in] graphExec - instance of executable graph to launch. + * @param [in] stream - instance of stream in which to launch executable graph. + * @returns #hipSuccess, #hipErrorInvalidValue + * + * @warning : This API is marked as beta, meaning, while this is feature complete, + * it is still open to changes and may have outstanding issues. + */ +hipError_t hipGraphUpload(hipGraphExec_t graphExec, hipStream_t stream); + /** * @brief Destroys an executable graph * From da2888b9c4772d149b32e80afacdb48b0a1dd0e5 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 30 Jun 2022 13:29:27 +0530 Subject: [PATCH 3/3] SWDEV-341956 - update HIP Dtest with correct coordinates (#2755) Change-Id: Iaccc4251a984e690d680ce884c6bfcc1093703af [ROCm/hip commit: 2667f3e22531d5857ce5f1ba44b40b1e31738c86] --- .../hip/tests/src/deviceLib/hipTestAtomicAdd.cpp | 2 +- .../tests/src/dynamicLoading/bit_extract_kernel.cpp | 4 ++-- .../src/dynamicLoading/complex_loading_behavior.cpp | 4 ++-- .../hip/tests/src/hostcall/hipHostcallFuncCall.cpp | 2 +- .../tests/src/hostcall/hipHostcallPrintThings.cpp | 4 ++-- projects/hip/tests/src/kernel/hipShflTests.cpp | 2 +- projects/hip/tests/src/printf/hipPrintfBasic.cpp | 12 ++++++------ .../hip/tests/src/printf/hipPrintfManyDevices.cpp | 2 +- projects/hip/tests/src/printf/hipPrintfManyWaves.cpp | 10 +++++----- projects/hip/tests/src/texture/hipBindTex2DPitch.cpp | 4 ++-- .../tests/src/texture/hipNormalizedFloatValueTex.cpp | 2 +- 11 files changed, 24 insertions(+), 24 deletions(-) diff --git a/projects/hip/tests/src/deviceLib/hipTestAtomicAdd.cpp b/projects/hip/tests/src/deviceLib/hipTestAtomicAdd.cpp index 9d03249e10..25f818cf8b 100644 --- a/projects/hip/tests/src/deviceLib/hipTestAtomicAdd.cpp +++ b/projects/hip/tests/src/deviceLib/hipTestAtomicAdd.cpp @@ -75,7 +75,7 @@ bool p_atomicNoRet = false; template __global__ void atomicnoret_manywaves(T* C_d) { - size_t tid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); + size_t tid = (blockIdx.x * blockDim.x + threadIdx.x); switch (tid % 9) { case 0: atomicAddNoRet(C_d, INCREMENT_VALUE); diff --git a/projects/hip/tests/src/dynamicLoading/bit_extract_kernel.cpp b/projects/hip/tests/src/dynamicLoading/bit_extract_kernel.cpp index 3a5141d817..2647a6c441 100644 --- a/projects/hip/tests/src/dynamicLoading/bit_extract_kernel.cpp +++ b/projects/hip/tests/src/dynamicLoading/bit_extract_kernel.cpp @@ -21,8 +21,8 @@ THE SOFTWARE. extern "C" __global__ void bit_extract_kernel(uint32_t* C_d, const uint32_t* A_d, size_t N) { - size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); - size_t stride = hipBlockDim_x * hipGridDim_x; + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; for (size_t i = offset; i < N; i += stride) { #ifdef __HIP_PLATFORM_AMD__ diff --git a/projects/hip/tests/src/dynamicLoading/complex_loading_behavior.cpp b/projects/hip/tests/src/dynamicLoading/complex_loading_behavior.cpp index f13f4b5c18..d83e98a882 100644 --- a/projects/hip/tests/src/dynamicLoading/complex_loading_behavior.cpp +++ b/projects/hip/tests/src/dynamicLoading/complex_loading_behavior.cpp @@ -37,8 +37,8 @@ THE SOFTWARE. #include "test_common.h" __global__ void vector_add(float* C, float* A, float* B, size_t N) { - size_t offset = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; - size_t stride = hipBlockDim_x * hipGridDim_x; + size_t offset = blockIdx.x * blockDim.x + threadIdx.x; + size_t stride = blockDim.x * gridDim.x; for (size_t i = offset; i < N; i += stride) { C[i] = A[i] + B[i]; } diff --git a/projects/hip/tests/src/hostcall/hipHostcallFuncCall.cpp b/projects/hip/tests/src/hostcall/hipHostcallFuncCall.cpp index fd9923e372..d7ffc04bbc 100644 --- a/projects/hip/tests/src/hostcall/hipHostcallFuncCall.cpp +++ b/projects/hip/tests/src/hostcall/hipHostcallFuncCall.cpp @@ -36,7 +36,7 @@ static void callee(uint64_t* output, uint64_t* input) { } __global__ void kernel(uint64_t fptr, uint64_t* retval0, uint64_t* retval1) { - uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + uint tid = threadIdx.x + blockIdx.x * blockDim.x; uint64_t arg0 = (uint64_t)fptr; uint64_t arg1 = tid; uint64_t arg2 = 42; diff --git a/projects/hip/tests/src/hostcall/hipHostcallPrintThings.cpp b/projects/hip/tests/src/hostcall/hipHostcallPrintThings.cpp index 9e53aeb9eb..02119079c5 100644 --- a/projects/hip/tests/src/hostcall/hipHostcallPrintThings.cpp +++ b/projects/hip/tests/src/hostcall/hipHostcallPrintThings.cpp @@ -43,7 +43,7 @@ void print_things_0(ulong* output, ulong* input) { } __global__ void kernel0(ulong fptr, ulong* retval) { - uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + uint tid = threadIdx.x + blockIdx.x * blockDim.x; ulong arg0 = fptr; const char* str = "(%lu -> %lu)\n"; @@ -77,7 +77,7 @@ void print_things_1(ulong* output, const ulong* input) { } __global__ void kernel1(ulong fptr, ulong name, ulong* retval) { - uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + uint tid = threadIdx.x + blockIdx.x * blockDim.x; ulong arg0 = fptr; ulong arg1 = name; ulong arg2 = tid; diff --git a/projects/hip/tests/src/kernel/hipShflTests.cpp b/projects/hip/tests/src/kernel/hipShflTests.cpp index 92205c4054..e1467aaa17 100644 --- a/projects/hip/tests/src/kernel/hipShflTests.cpp +++ b/projects/hip/tests/src/kernel/hipShflTests.cpp @@ -40,7 +40,7 @@ THE SOFTWARE. // Device (Kernel) function, it must be void template __global__ void matrixTranspose(T* out, T* in, const int width) { - int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; + int x = blockDim.x * blockIdx.x + threadIdx.x; T val = in[x]; for (int i = 0; i < width; i++) { for (int j = 0; j < width; j++) out[i * width + j] = __shfl(val, j * width + i); diff --git a/projects/hip/tests/src/printf/hipPrintfBasic.cpp b/projects/hip/tests/src/printf/hipPrintfBasic.cpp index 8ad578cecc..6169dd01e4 100644 --- a/projects/hip/tests/src/printf/hipPrintfBasic.cpp +++ b/projects/hip/tests/src/printf/hipPrintfBasic.cpp @@ -36,7 +36,7 @@ THE SOFTWARE. DECLARE_DATA(); __global__ void kernel_uniform0(int *retval) { - uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + uint tid = threadIdx.x + blockIdx.x * blockDim.x; retval[tid] = printf("Hello World\n"); // In Hip-Rocclr, printf returns number of characters printed. // In Cuda, printf returns the number of arguments parsed. } @@ -78,7 +78,7 @@ static void test_uniform0(int *retval, uint num_blocks, } __global__ void kernel_uniform1(int *retval) { - uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + uint tid = threadIdx.x + blockIdx.x * blockDim.x; retval[tid] = printf("Six times Eight is %d\n", 42); } @@ -119,7 +119,7 @@ static void test_uniform1(int *retval, uint num_blocks, } __global__ void kernel_divergent0(int *retval) { - uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + uint tid = threadIdx.x + blockIdx.x * blockDim.x; retval[tid] = printf("Thread ID: %d\n", tid); } @@ -171,7 +171,7 @@ static void test_divergent0(int *retval, uint num_blocks, } __global__ void kernel_divergent1(int *retval) { - uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + uint tid = threadIdx.x + blockIdx.x * blockDim.x; if (tid % 2) { retval[tid] = printf("Hello World\n"); } else { @@ -222,7 +222,7 @@ static void test_divergent1(int *retval, uint num_blocks, __global__ void kernel_series(int *retval) { DECLARE_DATA(); - const uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + const uint tid = threadIdx.x + blockIdx.x * blockDim.x; int result = 0; result += printf("%s\n", msg_long1); result += printf("%s\n", msg_short); @@ -269,7 +269,7 @@ static void test_series(int *retval, uint num_blocks, uint threads_per_block) { } __global__ void kernel_divergent_loop() { - const uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + const uint tid = threadIdx.x + blockIdx.x * blockDim.x; int result = 0; for (int i = 0; i <= tid; ++i) { diff --git a/projects/hip/tests/src/printf/hipPrintfManyDevices.cpp b/projects/hip/tests/src/printf/hipPrintfManyDevices.cpp index 2e7d99461f..004757cac3 100644 --- a/projects/hip/tests/src/printf/hipPrintfManyDevices.cpp +++ b/projects/hip/tests/src/printf/hipPrintfManyDevices.cpp @@ -34,7 +34,7 @@ DECLARE_DATA(); __global__ void print_things() { DECLARE_DATA(); - uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + uint tid = threadIdx.x + blockIdx.x * blockDim.x; const char *msg[] = {msg_short, msg_long1, msg_long2}; printf("%s\n", msg[tid % 3]); diff --git a/projects/hip/tests/src/printf/hipPrintfManyWaves.cpp b/projects/hip/tests/src/printf/hipPrintfManyWaves.cpp index 0cb2b16c24..3bcb7fb4c1 100644 --- a/projects/hip/tests/src/printf/hipPrintfManyWaves.cpp +++ b/projects/hip/tests/src/printf/hipPrintfManyWaves.cpp @@ -38,7 +38,7 @@ DECLARE_DATA(); __global__ void kernel_mixed0(int *retval) { DECLARE_DATA(); - uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + uint tid = threadIdx.x + blockIdx.x * blockDim.x; // Three strings passed as divergent values to the same hostcall. const char *msg; @@ -107,7 +107,7 @@ static void test_mixed0(int *retval, uint num_blocks, uint threads_per_block) { __global__ void kernel_mixed1(int *retval) { DECLARE_DATA(); - const uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + const uint tid = threadIdx.x + blockIdx.x * blockDim.x; // Three strings passed to divergent hostcalls. switch (tid % 3) { @@ -173,7 +173,7 @@ static void test_mixed1(int *retval, uint num_blocks, uint threads_per_block) { __global__ void kernel_mixed2(int *retval) { DECLARE_DATA(); - const uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + const uint tid = threadIdx.x + blockIdx.x * blockDim.x; // Three different strings. All workitems print all three, but // in different orders. @@ -230,7 +230,7 @@ static void test_mixed2(int *retval, uint num_blocks, uint threads_per_block) { __global__ void kernel_mixed3(int *retval) { DECLARE_DATA(); - const uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + const uint tid = threadIdx.x + blockIdx.x * blockDim.x; int result = 0; result += printf("%s\n", msg_long1); @@ -300,7 +300,7 @@ static void test_mixed3(int *retval, uint num_blocks, uint threads_per_block) { } __global__ void kernel_numbers() { - uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + uint tid = threadIdx.x + blockIdx.x * blockDim.x; for (uint i = 0; i != 7; ++i) { uint base = tid * 21 + i * 3; printf("%d %d %d\n", base, base + 1, base + 2); diff --git a/projects/hip/tests/src/texture/hipBindTex2DPitch.cpp b/projects/hip/tests/src/texture/hipBindTex2DPitch.cpp index 41c99a8a61..3fb8d21f8b 100644 --- a/projects/hip/tests/src/texture/hipBindTex2DPitch.cpp +++ b/projects/hip/tests/src/texture/hipBindTex2DPitch.cpp @@ -33,8 +33,8 @@ texture tex; // texture object is a kernel argument __global__ void texture2dCopyKernel( TYPE_t* dst) { #if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT - int x = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; - int y = hipThreadIdx_y + hipBlockIdx_y * hipBlockDim_y; + int x = threadIdx.x + blockIdx.x * blockDim.x; + int y = threadIdx.y + blockIdx.y * blockDim.y; if ( (x< SIZE_W) && (y< SIZE_H) ){ dst[SIZE_W*y+x] = tex2D(tex, x, y); } diff --git a/projects/hip/tests/src/texture/hipNormalizedFloatValueTex.cpp b/projects/hip/tests/src/texture/hipNormalizedFloatValueTex.cpp index 2430f09906..0447c6aae2 100644 --- a/projects/hip/tests/src/texture/hipNormalizedFloatValueTex.cpp +++ b/projects/hip/tests/src/texture/hipNormalizedFloatValueTex.cpp @@ -60,7 +60,7 @@ template __global__ void normalizedValTextureTest(unsigned int numElements, float* pDst) { #if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT - unsigned int elementID = hipThreadIdx_x; + unsigned int elementID = threadIdx.x; if(elementID >= numElements) return; float coord =(float) elementID/numElements;