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-tests commit: 693b89b3dd]

projects/hip-tests/samples/0_Intro/bit_extract/README.md

@@ -1,6 +1,6 @@
 # bit_extract
 
-Show an application written directly in HIP which uses platform-specific check on __HIP_PLATFORM_HCC__ to enable use of 
+Show an application written directly in HIP which uses platform-specific check on __HIP_PLATFORM_AMD__ to enable use of
 an instruction that only exists on the HCC platform.
 
 See related [blog](http://gpuopen.com/platform-aware-coding-inside-hip/) demonstrating platform specialization.

projects/hip-tests/samples/0_Intro/bit_extract/bit_extract.cpp

@@ -39,9 +39,9 @@ __global__ void bit_extract_kernel(uint32_t* C_d, const uint32_t* A_d, size_t N)
     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/samples/0_Intro/module_api/launchKernelHcc.cpp

@@ -26,7 +26,7 @@ THE SOFTWARE.
 #include <fstream>
 #include <vector>
 
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
 #include <hip/hip_ext.h>
 #endif
 

projects/hip-tests/samples/0_Intro/square/square.hipref.cpp

@@ -57,7 +57,7 @@ int main(int argc, char* argv[]) {
     hipDeviceProp_t props;
     CHECK(hipGetDeviceProperties(&props, device /*deviceID*/));
     printf("info: running on device %s\n", props.name);
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
     printf("info: architecture on AMD GPU device is: %s\n", props.gcnArchName);
 #endif
     printf("info: allocate host mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0);

projects/hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp

@@ -813,12 +813,12 @@ int main(int argc, char* argv[]) {
     CommandStream* cs;
 
     if (p_blockingSync) {
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
         printf("setting BlockingSync for AMD\n");
         setenv("HIP_BLOCKING_SYNC", "1", 1);
 
 #endif
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
         printf("setting cudaDeviceBlockingSync\n");
         HIPCHECK(hipSetDeviceFlags(cudaDeviceBlockingSync));
 #endif

projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchEnqueueRateMT.cpp

@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 #include <stdio.h>
 #include "hip/hip_runtime.h"
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
 #include "hip/hip_ext.h"
 #endif
 #include <iostream>

projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp

@@ -18,7 +18,7 @@ THE SOFTWARE.
 */
 
 #include "hip/hip_runtime.h"
-#ifdef __HIP_PLATFORM_HCC__
+#ifdef __HIP_PLATFORM_AMD__
 #include "hip/hip_ext.h"
 #endif
 #include <iostream>
@@ -109,7 +109,7 @@ int main() {
     /***********************************************************************************/
 
     //Timing directly the dispatch
-#if defined(__HIP_PLATFORM_HCC__) && GENERIC_GRID_LAUNCH == 1 && defined(__HCC__)
+#if defined(__HIP_PLATFORM_AMD__) && GENERIC_GRID_LAUNCH == 1 && defined(__HCC__)
     for (auto i = 0; i < TOTAL_RUN_COUNT; ++i) {
         hipExtLaunchKernelGGL((EmptyKernel), dim3(NUM_GROUPS), dim3(GROUP_SIZE), 0, stream0, start, stop, 0);
         hipEventSynchronize(stop);

projects/hip-tests/samples/1_Utils/hipInfo/README.md

@@ -3,4 +3,4 @@
 Simple tool that prints properties for each device (from hipGetDeviceProperties), and compiler info.
     Properties includes all of the architectural feature flags for each device.
 
-Also demonstrates how to use platform-specific compilation path (testing `__HIP_PLATFORM_NVCC__` or `__HIP_PLATFORM_HCC__`)
+Also demonstrates how to use platform-specific compilation path (testing `__HIP_PLATFORM_AMD__` or `__HIP_PLATFORM_NVIDIA__`)

projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp

@@ -171,7 +171,7 @@ void printDeviceProp(int deviceId) {
     cout << endl;
 
 
-#ifdef __HIP_PLATFORM_NVCC__
+#ifdef __HIP_PLATFORM_NVIDIA__
     // Limits:
     cout << endl;
     printLimit(w1, cudaLimitStackSize, "bytes/thread");