Remove hip-hcc codes: Part three

1.Rename include/hip/hcc_detail/ as include/hip/amd_detail/

2.Rename include/hip/nvcc_detail/ as include/hip/nvidia_detail/

3.Create __HIP_PLATFORM_AMD__ to replace __HIP_PLATFORM_HCC__

4.Create __HIP_PLATFORM_NVIDIA__ to replace __HIP_PLATFORM_NVCC__

After hcc_detail, nvcc_detail, __HIP_PLATFORM_HCC__ and __HIP_PLATFORM_NVCC__
have been removed from upstream, they will be removed from hip runtime.

Change-Id: I1ae457effd739d6c25bca203c1724b026be21fce


[ROCm/hip commit: c2adc70d4d]

projects/hip/tests/src/cppstd/is_callable_test.hpp

@@ -20,7 +20,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 
-#include <hip/hcc_detail/helpers.hpp>
+#include <hip/amd_detail/helpers.hpp>
 #include <memory>
 
 using hip_impl::is_callable;

projects/hip/tests/src/deviceLib/hipComplex.cpp

@@ -21,7 +21,7 @@ THE SOFTWARE.
 #include <iostream>
 #include "hip/hip_runtime.h"
 #include "hip/hip_runtime_api.h"
-#include "hip/hcc_detail/hip_complex.h"
+#include "hip/amd_detail/hip_complex.h"
 
 #define LEN 64
 #define SIZE 64 << 2

projects/hip/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp

@@ -102,7 +102,7 @@ __host__ void double_precision_math_functions() {
     nearbyint(0.0);
     // nextafter(0.0);
     fX = 1.0;  // norm(1, &fX);
-#if defined(__HIP_PLATFORM_HCC__)
+#if defined(__HIP_PLATFORM_AMD__)
     // norm3d(1.0, 0.0, 0.0);
     // norm4d(1.0, 0.0, 0.0, 0.0);
 #endif
@@ -113,11 +113,11 @@ __host__ void double_precision_math_functions() {
 
     remainder(2.0, 1.0);
     remquo(1.0, 2.0, &iX);
-#if defined(__HIP_PLATFORM_HCC__)
+#if defined(__HIP_PLATFORM_AMD__)
     // rhypot(0.0, 1.0);
 #endif
     rint(1.0);
-#if defined(__HIP_PLATFORM_HCC__)
+#if defined(__HIP_PLATFORM_AMD__)
     fX = 1.0;  // rnorm(1, &fX);
     // rnorm3d(0.0, 0.0, 1.0);
     // rnorm4d(0.0, 0.0, 0.0, 1.0);

projects/hip/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp

@@ -66,7 +66,7 @@ __host__ void single_precision_math_functions() {
     expm1f(0.0f);
     fabsf(1.0f);
     fdimf(1.0f, 0.0f);
-#if defined(__HIP_PLATFORM_HCC__)
+#if defined(__HIP_PLATFORM_AMD__)
     // fdividef(0.0f, 1.0f);
 #endif
     floorf(0.0f);
@@ -100,7 +100,7 @@ __host__ void single_precision_math_functions() {
     nanf("1");
     nearbyintf(0.0f);
     // nextafterf(0.0f);
-#if defined(__HIP_PLATFORM_HCC__)
+#if defined(__HIP_PLATFORM_AMD__)
     // norm3df(1.0f, 0.0f, 0.0f);
     // norm4df(1.0f, 0.0f, 0.0f, 0.0f);
 #endif
@@ -111,11 +111,11 @@ __host__ void single_precision_math_functions() {
     // rcbrtf(1.0f);
     remainderf(2.0f, 1.0f);
     remquof(1.0f, 2.0f, &iX);
-#if defined(__HIP_PLATFORM_HCC__)
+#if defined(__HIP_PLATFORM_AMD__)
     // rhypotf(0.0f, 1.0f);
 #endif
     rintf(1.0f);
-#if defined(__HIP_PLATFORM_HCC__)
+#if defined(__HIP_PLATFORM_AMD__)
     // rnorm3df(0.0f, 0.0f, 1.0f);
     // rnorm4df(0.0f, 0.0f, 0.0f, 1.0f);
     fX = 1.0f;  // rnormf(1, &fX);

projects/hip/tests/src/deviceLib/hip_ballot.cpp

@@ -34,7 +34,7 @@ __global__ void gpu_ballot(unsigned int* device_ballot, int Num_Warps_per_Block,
                            int pshift) {
     int tid = threadIdx.x + blockIdx.x * blockDim.x;
     const unsigned int warp_num = threadIdx.x >> pshift;
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
     atomicAdd(&device_ballot[warp_num + blockIdx.x * Num_Warps_per_Block],
               __popcll(__ballot(tid - 245)));
 #else

projects/hip/tests/src/deviceLib/hip_test_ldg.cpp

@@ -35,7 +35,7 @@ THE SOFTWARE.
 #include "hip/hip_vector_types.h"
 #include "test_common.h"
 
-#if (__hcc_workweek__ >= 16164) || defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_CLANG_ONLY__)
+#if (__hcc_workweek__ >= 16164) || defined(__HIP_PLATFORM_NVIDIA__) || defined(__HIP_CLANG_ONLY__)
 
 #define HIP_ASSERT(x) (assert((x) == hipSuccess))
 

projects/hip/tests/src/dynamicLoading/bit_extract_kernel.cpp

@@ -25,9 +25,9 @@ extern "C" __global__ void bit_extract_kernel(uint32_t* C_d, const uint32_t*
   size_t stride = hipBlockDim_x * hipGridDim_x;
 
   for (size_t i = offset; i < N; i += stride) {
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
     C_d[i] = __bitextract_u32(A_d[i], 8, 4);
-#else /* defined __HIP_PLATFORM_NVCC__ or other path */
+#else /* defined __HIP_PLATFORM_NVIDIA__ or other path */
     C_d[i] = ((A_d[i] & 0xf00) >> 8);
 #endif
   }

projects/hip/tests/src/experimental/xcompile/gapi.sh

@@ -1,9 +1,9 @@
 #!/bin/bash
 
-gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gHipApi.o
+gcc -D__HIP_PLATFORM_AMD__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gHipApi.o
 
-gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c hHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o hHipApi.o
+gcc -D__HIP_PLATFORM_AMD__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c hHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o hHipApi.o
 
-gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include hHipApi.o gHipApi.o ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -lm -o gApi
+gcc -D__HIP_PLATFORM_AMD__= -I${HIP_PATH}/include -I${HCC_HOME}/include hHipApi.o gHipApi.o ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -lm -o gApi
 
 

projects/hip/tests/src/experimental/xcompile/ghipapi.sh

@@ -1,8 +1,8 @@
 #!/bin/bash
 
-gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gHipApi.o
+gcc -D__HIP_PLATFORM_AMD__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gHipApi.o
 
-gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c hHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o hHipApi.o
+gcc -D__HIP_PLATFORM_AMD__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c hHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o hHipApi.o
 
 hipcc hHipApi.o gHipApi.o -o gHipApi
 

projects/hip/tests/src/experimental/xcompile/gxxapi.sh

@@ -1,8 +1,8 @@
 #!/bin/bash
 
-g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi.o
+g++ -D__HIP_PLATFORM_AMD__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi.o
 
-g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c hxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o hxxHipApi.o
+g++ -D__HIP_PLATFORM_AMD__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c hxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o hxxHipApi.o
 
-g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include gxxHipApi.o hxxHipApi.o ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxApi
+g++ -D__HIP_PLATFORM_AMD__= -I${HIP_PATH}/include -I${HCC_HOME}/include gxxHipApi.o hxxHipApi.o ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxApi
 

projects/hip/tests/src/experimental/xcompile/gxxhipapi.sh

@@ -1,8 +1,8 @@
 #!/bin/bash
 
-g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi.o
+g++ -D__HIP_PLATFORM_AMD__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi.o
 
 hipcc -c hxxHipApi.cpp -o hxxHipApi.o
 
-g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include gxxHipApi.o hxxHipApi.o ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi
+g++ -D__HIP_PLATFORM_AMD__= -I${HIP_PATH}/include -I${HCC_HOME}/include gxxHipApi.o hxxHipApi.o ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi
 

projects/hip/tests/src/experimental/xcompile/hipapig.sh

@@ -1,3 +1,3 @@
 #!/bin/bash
 
-gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include gApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -lm -o gApi
+gcc -D__HIP_PLATFORM_AMD__= -I${HIP_PATH}/include -I${HCC_HOME}/include gApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -lm -o gApi

projects/hip/tests/src/experimental/xcompile/hipapigxx.sh

@@ -1,3 +1,3 @@
 #!/bin/bash
 
-g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include gxxApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxApi
+g++ -D__HIP_PLATFORM_AMD__= -I${HIP_PATH}/include -I${HCC_HOME}/include gxxApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxApi

projects/hip/tests/src/experimental/xcompile/hipg.sh

@@ -1,6 +1,6 @@
 #!/bin/bash
 
-gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gHipApi.o
+gcc -D__HIP_PLATFORM_AMD__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gHipApi.o
 
 hipcc -c hHip.c -o hHip.o
 

projects/hip/tests/src/experimental/xcompile/hipgapi.sh

@@ -1,6 +1,6 @@
 #!/bin/bash
 
-gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gHipApi.o
+gcc -D__HIP_PLATFORM_AMD__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gHipApi.o
 
 hipcc -c hHipApi.c -o hHipApi.o
 

projects/hip/tests/src/experimental/xcompile/hipgxx.sh

@@ -1,6 +1,6 @@
 #!/bin/bash
 
-g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi.o
+g++ -D__HIP_PLATFORM_AMD__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi.o
 
 hipcc -c hxxHip.cpp -o hxxHip.o
 

projects/hip/tests/src/experimental/xcompile/hipgxxapi.sh

@@ -1,6 +1,6 @@
 #!/bin/bash
 
-g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi.o
+g++ -D__HIP_PLATFORM_AMD__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi.o
 
 hipcc -c hxxHipApi.cpp -o hxxHipApi.o
 

projects/hip/tests/src/g++/hipMalloc.cpp

@@ -18,7 +18,7 @@
  * */
 
 /* HIT_START
- * BUILD_CMD: hipMalloc %cxx -D__HIP_PLATFORM_HCC__ -I%hip-path/include -I%rocm-path/include %S/%s -Wl,--rpath=%hip-path/lib %hip-path/lib/libamdhip64.so -o %T/%t -std=c++11 EXCLUDE_HIP_PLATFORM nvidia EXCLUDE_HIP_LIB_TYPE static
+ * BUILD_CMD: hipMalloc %cxx -D__HIP_PLATFORM_AMD__ -I%hip-path/include -I/opt/rocm/include %S/%s -Wl,--rpath=%hip-path/lib %hip-path/lib/libamdhip64.so -o %T/%t -std=c++11 EXCLUDE_HIP_PLATFORM nvidia EXCLUDE_HIP_LIB_TYPE static
  * TEST: %t EXCLUDE_HIP_PLATFORM nvidia EXCLUDE_HIP_LIB_TYPE static
  * HIT_END
  */

projects/hip/tests/src/gcc/LaunchKernel.c

@@ -20,7 +20,7 @@
 
 /* HIT_START
  * BUILD_CMD: gpu.o %hc -I%hip-path/include -g -c %S/gpu.cpp -o %T/gpu.o EXCLUDE_HIP_PLATFORM nvidia EXCLUDE_HIP_RUNTIME rocclr
- * BUILD_CMD: launchkernel.o %hc -D__HIP_PLATFORM_HCC__ -g -I%hip-path/include -c %S/LaunchKernel.c -o %T/launchkernel.o EXCLUDE_HIP_PLATFORM nvidia EXCLUDE_HIP_RUNTIME rocclr
+ * BUILD_CMD: launchkernel.o %hc -D__HIP_PLATFORM_AMD__ -g -I%hip-path/include -c %S/LaunchKernel.c -o %T/launchkernel.o EXCLUDE_HIP_PLATFORM nvidia EXCLUDE_HIP_RUNTIME rocclr
  * BUILD_CMD: LaunchKernel %hc %T/launchkernel.o %T/gpu.o -g -Wl,--rpath=%hip-path/lib %hip-path/lib/libamdhip64.so -o %T/%t DEPENDS gpu.o launchkernel.o EXCLUDE_HIP_PLATFORM nvidia EXCLUDE_HIP_RUNTIME rocclr
  * TEST: %t EXCLUDE_HIP_PLATFORM nvidia EXCLUDE_HIP_RUNTIME rocclr
  * HIT_END

projects/hip/tests/src/gcc/hipMalloc.c

@@ -18,8 +18,8 @@
  * */
 
 /* HIT_START
- * BUILD_CMD: hipMalloc %cc -D__HIP_PLATFORM_NVCC__ -I%hip-path/include -I/usr/local/cuda/include %S/%s -o %T/hipMalloc_nv -L/usr/local/cuda/lib64 -lcudart EXCLUDE_HIP_PLATFORM amd
- * BUILD_CMD: hipMalloc %cc -D__HIP_PLATFORM_HCC__ -I%hip-path/include %S/%s -Wl,--rpath=%hip-path/lib %hip-path/lib/libamdhip64.so -o %T/hipMalloc_hcc EXCLUDE_HIP_PLATFORM nvidia EXCLUDE_HIP_RUNTIME rocclr
+ * BUILD_CMD: hipMalloc %cc -D__HIP_PLATFORM_NVIDIA__ -I%hip-path/include -I/usr/local/cuda/include %S/%s -o %T/hipMalloc_nv -L/usr/local/cuda/lib64 -lcudart EXCLUDE_HIP_PLATFORM amd
+ * BUILD_CMD: hipMalloc %cc -D__HIP_PLATFORM_AMD__ -I%hip-path/include %S/%s -Wl,--rpath=%hip-path/lib %hip-path/lib/libamdhip64.so -o %T/hipMalloc_hcc EXCLUDE_HIP_PLATFORM nvidia EXCLUDE_HIP_RUNTIME rocclr
  * TEST: hipMalloc_nv EXCLUDE_HIP_PLATFORM amd
  * TEST: hipMalloc_hcc EXCLUDE_HIP_PLATFORM nvidia EXCLUDE_HIP_RUNTIME rocclr
  * HIT_END

projects/hip/tests/src/hipPointerAttrib.cpp

@@ -27,7 +27,7 @@ THE SOFTWARE.
 #include "test_common.h"
 #include <vector>
 
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
 #include "hc_am.hpp"
 
 #endif
@@ -289,7 +289,7 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) {
         reference[i]._pointer = ptr;
     }
 
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
     if (p_verbose & 0x2) {
         printf("Tracker after insertions:\n");
         hc::am_memtracker_print();
@@ -327,7 +327,7 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) {
         }
     }
 
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
     if (p_verbose & 0x2) {
         printf("Tracker after cleanup:\n");
         hc::am_memtracker_print();
@@ -392,7 +392,7 @@ void thread_query(void* ptr, const hipPointerAttribute_t* refAttrib) {
 }
 
 
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
 //---
 // Add pointers to tracker very quickly, then remove them quickly:
 enum Dir { Up, Down };
@@ -511,7 +511,7 @@ int main(int argc, char* argv[]) {
     }
 
 
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
     if (p_tests & 0x10) {
         srand(0x400);
         testMultiThreaded_2();

projects/hip/tests/src/kernel/hipDynamicShared.cpp

@@ -49,7 +49,7 @@ template <typename T>
 __global__ void testExternSharedKernel(const T* A_d, const T* B_d, T* C_d,
                                        size_t numElements, size_t groupElements) {
     // declare dynamic shared memory
-#if defined(__HIP_PLATFORM_HCC__)
+#if defined(__HIP_PLATFORM_AMD__)
     HIP_DYNAMIC_SHARED(T, sdata)
 #else
     HIP_DYNAMIC_SHARED(__align__(sizeof(T)) unsigned char, my_sdata)

projects/hip/tests/src/kernel/inline_asm_vadd.cpp

@@ -36,11 +36,11 @@ OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWA
 __global__ void vadd_asm(float* out, float* in) {
     int i = blockDim.x * blockIdx.x + threadIdx.x;
 
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
     asm volatile("add.f32 %0,%1,%2;" : "=f"(out[i]) : "f"(in[i]), "f"(out[i]));
 #endif
 
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
     asm volatile("v_add_f32_e32 %0, %1, %2" : "=v"(out[i]) : "v"(in[i]), "v"(out[i]));
 #endif
 }

projects/hip/tests/src/kernel/launch_bounds.cpp

@@ -80,7 +80,7 @@ int main(int argc, char* argv[]) {
 
     hipLaunchKernel(myKern, dim3(blocks), dim3(p_blockSize), 0, 0, C_d, A_d, N, 0);
 
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
     cudaFuncAttributes attrib;
     cudaFuncGetAttributes(&attrib, myKern);
     printf("binaryVersion = %d\n", attrib.binaryVersion);

projects/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp

@@ -136,7 +136,7 @@ int main(int argc, char* argv[]) {
     CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeCooperativeLaunch, props.cooperativeLaunch));
     CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeCooperativeMultiDeviceLaunch, props.cooperativeMultiDeviceLaunch));
 
-#ifndef __HIP_PLATFORM_NVCC__
+#ifndef __HIP_PLATFORM_NVIDIA__
     CHECK(test_hipDeviceGetHdpAddress(deviceId, hipDeviceAttributeHdpMemFlushCntl, props.hdpMemFlushCntl));
     CHECK(test_hipDeviceGetHdpAddress(deviceId, hipDeviceAttributeHdpRegFlushCntl, props.hdpRegFlushCntl));
 #endif

projects/hip/tests/src/runtimeApi/event/hipEventElapsedTime.cpp

@@ -35,7 +35,7 @@ void NegativeTests(){
         hipEvent_t start,end;
         float tms = 1.0f;
         HIPASSERT(hipEventElapsedTime(nullptr,start,end) == hipErrorInvalidValue);
-#ifndef __HIP_PLATFORM_NVCC__
+#ifndef __HIP_PLATFORM_NVIDIA__
         // On NVCC platform API throws seg fault hence skipping
         HIPASSERT(hipEventElapsedTime(&tms,nullptr,end) == hipErrorInvalidHandle);
         HIPASSERT(hipEventElapsedTime(&tms,start,nullptr) == hipErrorInvalidHandle);

projects/hip/tests/src/runtimeApi/memory/hipHostMalloc.cpp

@@ -150,7 +150,7 @@ int main() {
         int numElements = 1024 * 16;
         size_t sizeBytes = numElements * sizeof(int);
 
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
         {
             // Stimulate error condition:
             int* A = &numElements;

projects/hip/tests/src/runtimeApi/memory/hipMallocManaged_MultiScenario.cpp

@@ -206,7 +206,7 @@ bool NegativeTestsMallocManaged(int NumDevices) {
     IfTestPassed = false;
   }
 
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
   // The flag hipMemAttachHost is currently not supported therefore
   // api should return "hipErrorInvalidValue" for now
   err = hipMallocManaged(&A, 1024, hipMemAttachHost);
@@ -215,7 +215,7 @@ bool NegativeTestsMallocManaged(int NumDevices) {
            hipGetErrorString(err));
     IfTestPassed = false;
   }
-#endif  // __HIP_PLATFORM_HCC__
+#endif  // __HIP_PLATFORM_AMD__
 
   err = hipMallocManaged(NULL, 0, 0);
   if (hipErrorInvalidValue != err) {

projects/hip/tests/src/runtimeApi/memory/hipMemcpy3D.cpp

@@ -51,7 +51,7 @@ void runTest(int width,int height,int depth, hipChannelFormatKind formatKind)
     myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
     myparms.dstArray = arr;
     myparms.extent = make_hipExtent(width , height, depth);
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
     myparms.kind = cudaMemcpyHostToDevice;
 #else
     myparms.kind = hipMemcpyHostToDevice;
@@ -65,7 +65,7 @@ void runTest(int width,int height,int depth, hipChannelFormatKind formatKind)
     myparms.srcArray = arr;
     myparms.dstArray = arr1;
     myparms.extent = make_hipExtent(width, height, depth);
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
     myparms.kind = cudaMemcpyDeviceToDevice;
 #else
     myparms.kind = hipMemcpyDeviceToDevice;
@@ -82,7 +82,7 @@ void runTest(int width,int height,int depth, hipChannelFormatKind formatKind)
     myparms.dstPtr = make_hipPitchedPtr(hOutputData, width * sizeof(T), width, height);
     myparms.srcArray = arr1;
     myparms.extent = make_hipExtent(width, height, depth);
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
     myparms.kind = cudaMemcpyDeviceToHost;
 #else
     myparms.kind = hipMemcpyDeviceToHost;

projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp

@@ -60,7 +60,7 @@ bool testhipMemset3D(int memsetval,int p_gpuDevice)
     myparms.dstPtr = make_hipPitchedPtr(A_h, width , numW, numH);
     myparms.srcPtr = devPitchedPtr;
     myparms.extent = extent;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
     myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
 #else
     myparms.kind = hipMemcpyDeviceToHost;
@@ -111,7 +111,7 @@ bool testhipMemset3DAsync(int memsetval,int p_gpuDevice)
     myparms.dstPtr = make_hipPitchedPtr(A_h, width , numW, numH);
     myparms.srcPtr = devPitchedPtr;
     myparms.extent = extent;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
     myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
 #else
     myparms.kind = hipMemcpyDeviceToHost;

projects/hip/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp

@@ -34,7 +34,7 @@ THE SOFTWARE.
 
 #define USE_HCC_MEMTRACKER 0 /* Debug flag to show the memtracker periodically */
 
-#if defined(__HIP_PLATFORM_HCC__) && !defined(__HIP_ROCclr__)
+#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_ROCclr__)
 #include <hc_am.hpp>
 #else
 #define USE_HCC_MEMTRACKER 0

projects/hip/tests/src/runtimeApi/module/hipModuleLoadUnloadStress.cpp

@@ -98,7 +98,7 @@ void testhipModuleLoadDataExUnloadStress() {
 
 int main(int argc, char* argv[]) {
   HipTest::parseStandardArguments(argc, argv, true);
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipInit(0));
   hipDevice_t device;
   hipCtx_t context;
@@ -112,7 +112,7 @@ int main(int argc, char* argv[]) {
   } else if (p_tests == 0x3) {
     testhipModuleLoadDataExUnloadStress();
   }
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   passed();

projects/hip/tests/src/runtimeApi/module/hipModuleNegative.cpp

@@ -101,7 +101,7 @@ void createRandomFile(const char* filename) {
 /**
  * Internal Function
  */
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
 void initHipCtx(hipCtx_t *pcontext) {
   HIPCHECK(hipInit(0));
   hipDevice_t device;
@@ -117,7 +117,7 @@ void initHipCtx(hipCtx_t *pcontext) {
 bool testhipModuleLoadNeg10() {
   bool TestPassed = false;
   hipError_t ret;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -127,7 +127,7 @@ bool testhipModuleLoadNeg10() {
     printf("Test Passed: Error Code Returned: '%s'(%d)\n",
            hipGetErrorString(ret), ret);
   }
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   return TestPassed;
@@ -141,7 +141,7 @@ bool testhipModuleLoadNeg11() {
   bool TestPassed = false;
   hipError_t ret;
   hipModule_t Module;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -151,7 +151,7 @@ bool testhipModuleLoadNeg11() {
     printf("Test Passed: Error Code Returned: '%s'(%d)\n",
            hipGetErrorString(ret), ret);
   }
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   return TestPassed;
@@ -168,7 +168,7 @@ bool testhipModuleLoadNeg12() {
   std::fstream fs;
   fs.open(FILENAME_EMPTY, std::ios::out);
   fs.close();
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -178,7 +178,7 @@ bool testhipModuleLoadNeg12() {
     printf("Test Passed: Error Code Returned: '%s'(%d)\n",
            hipGetErrorString(ret), ret);
   }
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   remove(FILENAME_EMPTY);
@@ -195,7 +195,7 @@ bool testhipModuleLoadNeg13() {
   hipModule_t Module;
   // Create a binary file with random numbers
   createRandomFile(FILENAME_RAND);
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -205,7 +205,7 @@ bool testhipModuleLoadNeg13() {
     printf("Test Passed: Error Code Returned: '%s'(%d)\n",
            hipGetErrorString(ret), ret);
   }
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   remove(FILENAME_RAND);
@@ -220,7 +220,7 @@ bool testhipModuleLoadNeg14() {
   bool TestPassed = false;
   hipError_t ret;
   hipModule_t Module;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -230,7 +230,7 @@ bool testhipModuleLoadNeg14() {
     printf("Test Passed: Error Code Returned: '%s'(%d)\n",
            hipGetErrorString(ret), ret);
   }
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   return TestPassed;
@@ -244,7 +244,7 @@ bool testhipModuleLoadNeg15() {
   bool TestPassed = false;
   hipError_t ret;
   hipModule_t Module;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -253,7 +253,7 @@ bool testhipModuleLoadNeg15() {
     printf("Test Passed: Error Code Returned: '%s'(%d)\n",
            hipGetErrorString(ret), ret);
   }
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   return TestPassed;
@@ -267,7 +267,7 @@ bool testhipModuleLoadDataNeg20() {
   bool TestPassed = false;
   hipError_t ret;
   auto buffer = load_file(CODEOBJ_FILE);
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -277,7 +277,7 @@ bool testhipModuleLoadDataNeg20() {
     printf("Test Passed: Error Code Returned: '%s'(%d)\n",
            hipGetErrorString(ret), ret);
   }
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   return TestPassed;
@@ -291,7 +291,7 @@ bool testhipModuleLoadDataNeg21() {
   bool TestPassed = false;
   hipError_t ret;
   hipModule_t Module;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -301,7 +301,7 @@ bool testhipModuleLoadDataNeg21() {
     printf("Test Passed: Error Code Returned: '%s'(%d)\n",
            hipGetErrorString(ret), ret);
   }
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   return TestPassed;
@@ -319,7 +319,7 @@ bool testhipModuleLoadDataNeg22() {
   createRandomFile(FILENAME_RAND);
   // Open the code object file and copy it in a buffer
   auto buffer = load_file(FILENAME_RAND);
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -329,7 +329,7 @@ bool testhipModuleLoadDataNeg22() {
     printf("Test Passed: Error Code Returned: '%s'(%d)\n",
            hipGetErrorString(ret), ret);
   }
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   remove(FILENAME_RAND);
@@ -344,7 +344,7 @@ bool testhipModuleLoadDataExNeg30() {
   hipError_t ret;
   // Open the code object file and copy it in a buffer
   auto buffer = load_file(CODEOBJ_FILE);
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -354,7 +354,7 @@ bool testhipModuleLoadDataExNeg30() {
     printf("Test Passed: Error Code Returned: '%s'(%d)\n",
            hipGetErrorString(ret), ret);
   }
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   return TestPassed;
@@ -368,7 +368,7 @@ bool testhipModuleLoadDataExNeg31() {
   bool TestPassed = false;
   hipError_t ret;
   hipModule_t Module;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -378,7 +378,7 @@ bool testhipModuleLoadDataExNeg31() {
     printf("Test Passed: Error Code Returned: '%s'(%d)\n",
            hipGetErrorString(ret), ret);
   }
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   return TestPassed;
@@ -396,7 +396,7 @@ bool testhipModuleLoadDataExNeg32() {
   createRandomFile(FILENAME_RAND);
   // Open the code object file and copy it in a buffer
   auto buffer = load_file(FILENAME_RAND);
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -406,7 +406,7 @@ bool testhipModuleLoadDataExNeg32() {
     printf("Test Passed: Error Code Returned: '%s'(%d)\n",
            hipGetErrorString(ret), ret);
   }
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   remove(FILENAME_RAND);
@@ -421,7 +421,7 @@ bool testhipModuleGetFunctionNeg40() {
   bool TestPassed = false;
   hipError_t ret;
   hipModule_t Module;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -433,7 +433,7 @@ bool testhipModuleGetFunctionNeg40() {
            hipGetErrorString(ret), ret);
   }
   HIPCHECK(hipModuleUnload(Module));
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   return TestPassed;
@@ -448,7 +448,7 @@ bool testhipModuleGetFunctionNeg41() {
   hipError_t ret;
   hipModule_t Module;
   hipFunction_t Function;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -458,7 +458,7 @@ bool testhipModuleGetFunctionNeg41() {
     printf("Test Passed: Error Code Returned: '%s'(%d)\n",
            hipGetErrorString(ret), ret);
   }
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   return TestPassed;
@@ -473,7 +473,7 @@ bool testhipModuleGetFunctionNeg42() {
   hipError_t ret;
   hipModule_t Module;
   hipFunction_t Function;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -485,7 +485,7 @@ bool testhipModuleGetFunctionNeg42() {
            hipGetErrorString(ret), ret);
   }
   HIPCHECK(hipModuleUnload(Module));
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   return TestPassed;
@@ -500,7 +500,7 @@ bool testhipModuleGetFunctionNeg43() {
   hipError_t ret;
   hipModule_t Module;
   hipFunction_t Function;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -512,7 +512,7 @@ bool testhipModuleGetFunctionNeg43() {
            hipGetErrorString(ret), ret);
   }
   HIPCHECK(hipModuleUnload(Module));
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   return TestPassed;
@@ -527,7 +527,7 @@ bool testhipModuleGetFunctionNeg44() {
   hipError_t ret;
   hipModule_t Module;
   hipFunction_t Function;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -539,7 +539,7 @@ bool testhipModuleGetFunctionNeg44() {
     printf("Test Passed: Error Code Returned: '%s'(%d)\n",
            hipGetErrorString(ret), ret);
   }
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   return TestPassed;
@@ -554,7 +554,7 @@ bool testhipModuleGetFunctionNeg45() {
   hipError_t ret;
   hipModule_t Module;
   hipFunction_t Function;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -566,7 +566,7 @@ bool testhipModuleGetFunctionNeg45() {
            hipGetErrorString(ret), ret);
   }
   HIPCHECK(hipModuleUnload(Module));
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   return TestPassed;
@@ -581,7 +581,7 @@ bool testhipModuleGetGlobalNeg50() {
   hipError_t ret;
   hipModule_t Module;
   size_t deviceGlobalSize;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -593,7 +593,7 @@ bool testhipModuleGetGlobalNeg50() {
            hipGetErrorString(ret), ret);
   }
   HIPCHECK(hipModuleUnload(Module));
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   return TestPassed;
@@ -608,7 +608,7 @@ bool testhipModuleGetGlobalNeg51() {
   hipError_t ret;
   hipModule_t Module;
   hipDeviceptr_t deviceGlobal;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -620,7 +620,7 @@ bool testhipModuleGetGlobalNeg51() {
            hipGetErrorString(ret), ret);
   }
   HIPCHECK(hipModuleUnload(Module));
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   return TestPassed;
@@ -636,7 +636,7 @@ bool testhipModuleGetGlobalNeg52() {
   hipModule_t Module;
   hipDeviceptr_t deviceGlobal;
   size_t deviceGlobalSize;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -648,7 +648,7 @@ bool testhipModuleGetGlobalNeg52() {
            hipGetErrorString(ret), ret);
   }
   HIPCHECK(hipModuleUnload(Module));
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   return TestPassed;
@@ -664,7 +664,7 @@ bool testhipModuleGetGlobalNeg53() {
   hipModule_t Module;
   hipDeviceptr_t deviceGlobal;
   size_t deviceGlobalSize;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -676,7 +676,7 @@ bool testhipModuleGetGlobalNeg53() {
            hipGetErrorString(ret), ret);
   }
   HIPCHECK(hipModuleUnload(Module));
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   return TestPassed;
@@ -692,7 +692,7 @@ bool testhipModuleGetGlobalNeg54() {
   hipModule_t Module;
   hipDeviceptr_t deviceGlobal;
   size_t deviceGlobalSize;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -704,7 +704,7 @@ bool testhipModuleGetGlobalNeg54() {
            hipGetErrorString(ret), ret);
   }
   HIPCHECK(hipModuleUnload(Module));
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   return TestPassed;
@@ -720,7 +720,7 @@ bool testhipModuleGetGlobalNeg55() {
   hipModule_t Module;
   hipDeviceptr_t deviceGlobal;
   size_t deviceGlobalSize;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -732,7 +732,7 @@ bool testhipModuleGetGlobalNeg55() {
     printf("Test Passed: Error Code Returned: '%s'(%d)\n",
            hipGetErrorString(ret), ret);
   }
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   return TestPassed;
@@ -748,7 +748,7 @@ bool testhipModuleGetGlobalNeg56() {
   hipModule_t Module;
   hipDeviceptr_t deviceGlobal;
   size_t deviceGlobalSize;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -758,7 +758,7 @@ bool testhipModuleGetGlobalNeg56() {
     printf("Test Passed: Error Code Returned: '%s'(%d)\n",
            hipGetErrorString(ret), ret);
   }
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   return TestPassed;
@@ -773,7 +773,7 @@ bool testhipModuleLoadNeg60() {
   bool TestPassed = true;
   hipError_t ret;
   hipModule_t Module;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   hipCtx_t context;
   initHipCtx(&context);
 #endif
@@ -793,7 +793,7 @@ bool testhipModuleLoadNeg60() {
   } else {
     TestPassed &= false;
   }
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipCtxDestroy(context));
 #endif
   return TestPassed;

projects/hip/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp

@@ -59,7 +59,7 @@ THE SOFTWARE.
 #define GRIDDIMZ 1
 #define BLOCKDIMZ 1
 
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
 
 #define CTX_CREATE() \
   hipCtx_t context;\
@@ -79,14 +79,14 @@ void initHipCtx(hipCtx_t *pcontext) {
   HIPCHECK(hipCtxCreate(pcontext, 0, device));
 }
 
-#else  // __HIP_PLATFORM_NVCC__
+#else  // __HIP_PLATFORM_NVIDIA__
 
 #define CTX_CREATE()
 #define CTX_DESTROY()
 #define ARRAY_DESTROY(array) HIPCHECK(hipFreeArray(array));
 #define HIP_TEX_REFERENCE textureReference*
 #define HIP_ARRAY hipArray*
-#endif  // __HIP_PLATFORM_NVCC__
+#endif  // __HIP_PLATFORM_NVIDIA__
 
 std::atomic<int> g_thTestPassed(1);
 
@@ -271,7 +271,7 @@ template <class T, class T1> void copyBuffer2Array(unsigned int width,
                                                    ) {
   hip_Memcpy2D copyParam;
   memset(&copyParam, 0, sizeof(copyParam));
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   copyParam.dstMemoryType = CU_MEMORYTYPE_ARRAY;
   copyParam.srcMemoryType = CU_MEMORYTYPE_HOST;
   copyParam.dstArray = *array;
@@ -293,7 +293,7 @@ template <class T> void assignArray2TexRef(hipArray_Format format,
                                            T array
                                            ) {
   HIP_TEX_REFERENCE texref;
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   HIPCHECK(hipModuleGetTexRef(&texref, Module, texRefName));
   HIPCHECK(hipTexRefSetAddressMode(texref, 0, CU_TR_ADDRESS_MODE_WRAP));
   HIPCHECK(hipTexRefSetAddressMode(texref, 1, CU_TR_ADDRESS_MODE_WRAP));
@@ -351,7 +351,7 @@ template <class T> bool testTexType(hipArray_Format format,
   HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
   HIP_ARRAY array;
   allocInitArray(width, height, format, &array);
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   // Copy from hData to array using hipMemcpyParam2D
   copyBuffer2Array<T, HIP_ARRAY*>(width, height, hData, &array);
   // Get tex reference from the loaded kernel file
@@ -438,7 +438,7 @@ template <class T> bool testTexMultStream(const std::vector<char>& buffer,
   HIPCHECK(hipModuleLoadData(&Module, &buffer[0]));
   HIP_ARRAY array;
   allocInitArray(width, height, format, &array);
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
   // Copy from hData to array using hipMemcpyParam2D
   copyBuffer2Array<T, HIP_ARRAY*>(width, height, hData, &array);
   // Get tex reference from the loaded kernel file

projects/hip/tests/src/runtimeApi/occupancy/hipOccupancyMaxPotentialBlockSize.cpp

@@ -79,7 +79,7 @@ bool argValidation() {
     TestPassed &= false;
   }
 
-#ifndef __HIP_PLATFORM_NVCC__
+#ifndef __HIP_PLATFORM_NVIDIA__
   // nvcc doesnt support kernelfunc(NULL) for api
   ret = hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, NULL, 0, 0);
   if (ret != hipErrorInvalidValue) {

projects/hip/tests/src/runtimeApi/stream/hipStreamACb_AltEnqueue.cpp

@@ -37,7 +37,7 @@ THE SOFTWARE.
 #include "test_common.h"
 
 
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
 #define HIPRT_CB
 #endif
 

projects/hip/tests/src/runtimeApi/stream/hipStreamACb_MStrm_Mgpu.cpp

@@ -33,7 +33,7 @@
 #include "hip/hip_runtime.h"
 #include "test_common.h"
 
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
 #define HIPRT_CB
 #endif
 

projects/hip/tests/src/runtimeApi/stream/hipStreamACb_MultiCalls.cpp

@@ -34,7 +34,7 @@
 #include "hip/hip_runtime.h"
 #include "test_common.h"
 
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
 #define HIPRT_CB
 #endif
 

projects/hip/tests/src/runtimeApi/stream/hipStreamACb_MultiThread.cpp

@@ -36,7 +36,7 @@ THE SOFTWARE.
 #include "hip/hip_runtime.h"
 #include "test_common.h"
 
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
 #define HIPRT_CB
 #endif
 

projects/hip/tests/src/runtimeApi/stream/hipStreamACb_StrmSyncTiming.cpp

@@ -34,7 +34,7 @@
 #include "hip/hip_runtime.h"
 #include "test_common.h"
 
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
 #define HIPRT_CB
 #endif
 

projects/hip/tests/src/runtimeApi/stream/hipStreamACb_ThrdBehaviour.cpp

@@ -32,7 +32,7 @@
 #include "hip/hip_runtime.h"
 #include "test_common.h"
 
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
 #define HIPRT_CB
 #endif
 

projects/hip/tests/src/runtimeApi/stream/hipStreamACb_order.cpp

@@ -31,7 +31,7 @@
 #include <atomic>
 #include "hip/hip_runtime.h"
 #include "test_common.h"
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
 #define HIPRT_CB
 #endif
 

projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp

@@ -29,7 +29,7 @@ THE SOFTWARE.
 #include "hip/hip_runtime.h"
 #include "test_common.h"
 
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
 #define HIPRT_CB
 #endif
 

projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackCatch.cpp

@@ -30,7 +30,7 @@ using TestQueues = std::tuple<QueueHipRt<true>, QueueHipRt<false>>;
 #define HIP_ASSERT(x) (assert((x)==hipSuccess))
 #define HIP_ASSERT_IGNORE(x,ign) auto err=x; HIP_ASSERT(err==ign ? hipSuccess : err)
 
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
   #define HIPRT_CB
 #endif
 

projects/hip/tests/src/test_common.h

@@ -379,7 +379,7 @@ void freeArrays(T* A_d, T* B_d, T* C_d, T* A_h, T* B_h, T* C_h, bool usePinnedHo
     freeArraysForHost(A_h, B_h, C_h, usePinnedHost);
 }
 
-#if defined(__HIP_PLATFORM_HCC__)
+#if defined(__HIP_PLATFORM_AMD__)
 template <typename T>
 void initArrays2DPitch(T** A_d, T** B_d, T** C_d, size_t* pitch_A, size_t* pitch_B, size_t* pitch_C,
                        size_t numW, size_t numH) {

projects/hip/tests/src/texture/hipNormalizedFloatValueTex.cpp

@@ -116,7 +116,7 @@ int main(int argc, char** argv)
     hipDeviceProp_t props;
     HIPCHECK(hipGetDeviceProperties(&props, device));
     std::cout << "Device :: " << props.name << std::endl;
-    #ifdef __HIP_PLATFORM_HCC__
+    #ifdef __HIP_PLATFORM_AMD__
     std::cout << "Arch - AMD GPU :: " << props.gcnArch << std::endl;
     #endif