AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-299127 - Merge 'develop' into 'amd-staging'

Browse Source 
Change-Id: I3025b602e1ae16520f26b1f9bde2190625e28704


[ROCm/hip commit: 3483144a68]
This commit is contained in:
Jenkins
2022-06-30 19:10:25 -04:00
parent 6f1857e13a da2888b9c4
commit ead03d1342
12 changed files with 48 additions and 35 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/include/hip/hip_runtime_api.h
+24 -11
View File
@@ -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;
@@ -5617,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
*
projects/hip/tests/src/deviceLib/hipTestAtomicAdd.cpp
+1 -1
View File
@@ -75,7 +75,7 @@ bool p_atomicNoRet = false;
template <typename T>
__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);
projects/hip/tests/src/dynamicLoading/bit_extract_kernel.cpp
+2 -2
View File
@@ -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__
projects/hip/tests/src/dynamicLoading/complex_loading_behavior.cpp
+2 -2
View File
@@ -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];
}
projects/hip/tests/src/hostcall/hipHostcallFuncCall.cpp
+1 -1
View File
@@ -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;
projects/hip/tests/src/hostcall/hipHostcallPrintThings.cpp
+2 -2
View File
@@ -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;
projects/hip/tests/src/kernel/hipShflTests.cpp
+1 -1
View File
@@ -40,7 +40,7 @@ THE SOFTWARE.
// Device (Kernel) function, it must be void
template <typename T>
__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);
projects/hip/tests/src/printf/hipPrintfBasic.cpp
+6 -6
View File
@@ -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) {
projects/hip/tests/src/printf/hipPrintfManyDevices.cpp
+1 -1
View File
@@ -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]);
projects/hip/tests/src/printf/hipPrintfManyWaves.cpp
+5 -5
View File
@@ -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);
projects/hip/tests/src/texture/hipBindTex2DPitch.cpp
+2 -2
View File
@@ -33,8 +33,8 @@ texture<TYPE_t, 2, hipReadModeElementType> 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);
}
projects/hip/tests/src/texture/hipNormalizedFloatValueTex.cpp
+1 -1
View File
@@ -60,7 +60,7 @@ template<typename T>
__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;