diff --git a/projects/hip-tests/samples/0_Intro/bit_extract/README.md b/projects/hip-tests/samples/0_Intro/bit_extract/README.md index 76f8906595..69995721b0 100644 --- a/projects/hip-tests/samples/0_Intro/bit_extract/README.md +++ b/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. diff --git a/projects/hip-tests/samples/0_Intro/bit_extract/bit_extract.cpp b/projects/hip-tests/samples/0_Intro/bit_extract/bit_extract.cpp index 1f7770ccb5..dad27f759d 100644 --- a/projects/hip-tests/samples/0_Intro/bit_extract/bit_extract.cpp +++ b/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 } diff --git a/projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp b/projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp index 90e569c5bc..5bb1bd3e96 100644 --- a/projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp +++ b/projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp @@ -26,7 +26,7 @@ THE SOFTWARE. #include #include -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ #include #endif diff --git a/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp b/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp index a2b21f17e4..300a875997 100644 --- a/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp +++ b/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); diff --git a/projects/hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp b/projects/hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp index 21b5505623..37eb0845b1 100644 --- a/projects/hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp +++ b/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 diff --git a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchEnqueueRateMT.cpp b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchEnqueueRateMT.cpp index 8b58dd55e6..81cae36296 100644 --- a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchEnqueueRateMT.cpp +++ b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchEnqueueRateMT.cpp @@ -19,7 +19,7 @@ THE SOFTWARE. #include #include "hip/hip_runtime.h" -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ #include "hip/hip_ext.h" #endif #include diff --git a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp index b528b0c75d..fb4ab8455c 100644 --- a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp +++ b/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 @@ -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); diff --git a/projects/hip-tests/samples/1_Utils/hipInfo/README.md b/projects/hip-tests/samples/1_Utils/hipInfo/README.md index 1c2749880a..d52f8d8fef 100644 --- a/projects/hip-tests/samples/1_Utils/hipInfo/README.md +++ b/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__`) diff --git a/projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp b/projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp index e5998bfb64..e15ce583b9 100644 --- a/projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp +++ b/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");