Merge pull request #22 from oneapi-src/tsne/code_update/2309

[tsne] code update + cmake changes

Author: mgrabban

tsne/CUDA/CMakeLists.txt

@@ -38,12 +38,25 @@ set(CMAKE_CXX_EXTENSIONS OFF)
 
 option(USE_SM "Build for specific SM" OFF)
 
-# CMAKE_CXX_FLAGS
-if("${CMAKE_CXX_FLAGS}" STREQUAL "")
+set(DEF_WL_CXX_FLAGS           " ")
+set(DEF_GENERAL_CXX_FLAGS      " -O3 -ffast-math ")
+set(DEF_COMBINED_CXX_FLAGS     "${DEF_GENERAL_CXX_FLAGS} ${DEF_WL_CXX_FLAGS}")
+
+# -DCMAKE_CXX_FLAGS=" -blah -blah " overrides the default flags (BOTH general and WL specific)
+# -DOVERRIDE_GENERAL_CXX_FLAGS=" -blah -blah " overrides the general flags only (and not the workload specific flags)
+# passing in both CMAKE_CXX_FLAGS and OVERRIDE_GENERAL_CXX_FLAGS is not allowed, in order to prevent ambiguity
+
+if(NOT "${CMAKE_CXX_FLAGS}" STREQUAL "" AND NOT "${OVERRIDE_GENERAL_CXX_FLAGS}" STREQUAL "")
+    message(FATAL_ERROR "Both  CMAKE_CXX_FLAGS and OVERRIDE_GENERAL_CXX_FLAGS cannot be passed in together")
+elseif("${CMAKE_CXX_FLAGS}" STREQUAL "" AND "${OVERRIDE_GENERAL_CXX_FLAGS}" STREQUAL "")
     message(STATUS "Using DEFAULT compilation flags")
-    set(CMAKE_CXX_FLAGS ${CMAKE_CXX_FLAGS} " -O3 -ffast-math ")
-else()
-    message(STATUS "OVERRIDING DEFAULT compilation flags")
+    set(CMAKE_CXX_FLAGS "${DEF_COMBINED_CXX_FLAGS}")
+elseif(NOT "${OVERRIDE_GENERAL_CXX_FLAGS}" STREQUAL "")
+    message(STATUS "OVERRIDING GENERAL compilation flags")
+    set(CMAKE_CXX_FLAGS "${OVERRIDE_GENERAL_CXX_FLAGS}")
+    string(APPEND CMAKE_CXX_FLAGS ${DEF_WL_CXX_FLAGS})
+elseif(NOT "${CMAKE_CXX_FLAGS}" STREQUAL "")
+    message(STATUS "OVERRIDING GENERAL and WORKLOAD SPECIFIC compilation flags")
 endif()
 
 # CUDA_NVCC_FLAGS

tsne/CUDA/src/kernels/nbodyfft.cu

@@ -300,6 +300,7 @@ void DFT2D1gpu(float* din, thrust::complex<float>* dout, int num_rows, int num_c
     thrust::complex<float> sum, twiddle;
     angle = -2.0f * PI * fdividef((float)i, (float)num_cols);
     sum = 0.0f;
+#pragma unroll
     for (int k = 0; k < num_cols; ++k) {
         // sincosf(angle * k, &sinf, &cosf);
         // twiddle = thrust::complex<float>(cosf, sinf);
@@ -324,6 +325,7 @@ void DFT2D2gpu(thrust::complex<float>* din, thrust::complex<float>* dout, int nu
     thrust::complex<float> sum, twiddle;
     angle = -2.0f * PI * fdividef((float)i, (float)num_cols);
     sum = 0.0f;
+#pragma unroll
     for (int k = 0; k < num_cols; ++k) {
         // sincosf(angle * k, &sinf, &cosf);
         // twiddle = thrust::complex<float>(cosf, sinf);
@@ -348,6 +350,7 @@ void iDFT2D1gpu(thrust::complex<float>* din, thrust::complex<float>* dout, int n
     thrust::complex<float> sum, twiddle;
     angle = 2.0f * PI * fdividef((float)i, (float)num_cols);
     sum = 0.0f;
+#pragma unroll
     for (int k = 0; k < num_cols; ++k) {
         // sincosf(angle * k, &sinf, &cosf);
         // twiddle = thrust::complex<float>(cosf, sinf);
@@ -376,6 +379,7 @@ void iDFT2D2gpu(thrust::complex<float>* din, float* dout, int num_rows, int num_
     thrust::complex<float> twiddle;
     angle = 2.0f * PI * fdividef((float)i, (float)num_cols);
     sum = 0.0f;
+#pragma unroll
     for (int k = 0; k < num_cols; ++k) {
         // sincosf(angle * k, &sinf, &cosf);
         // twiddle = thrust::complex<float>(cosf, sinf);

tsne/CUDA/src/kernels/perplexity_search.cu

@@ -153,12 +153,12 @@ void tsnecuda::SearchPerplexity(
 
         // compute entropy of current row
         row_sum = tsnecuda::utils::ReduceSum(handle, pij, num_neighbors, num_points, 0);
-        GpuErrorCheck(cudaDeviceSynchronize());
+        // GpuErrorCheck(cudaDeviceSynchronize());
 
         // compute negative entropy
         thrust::transform(pij.begin(), pij.end(), entropy.begin(), tsnecuda::utils::FunctionalEntropy());
         neg_entropy = tsnecuda::utils::ReduceAlpha(handle, entropy, num_neighbors, num_points, -1.0f, 0);
-        GpuErrorCheck(cudaDeviceSynchronize());
+        // GpuErrorCheck(cudaDeviceSynchronize());
 
         // binary search for beta
         PerplexitySearchKernel<<<NBLOCKS2, BLOCKSIZE2>>>(

tsne/CUDA/src/utils/reduce_utils.cu

@@ -57,6 +57,7 @@ thrust::device_vector<float> tsnecuda::utils::ReduceAlpha(
                 thrust::raw_pointer_cast(d_matrix.data()), N,
                 thrust::raw_pointer_cast(ones.data()), 1, &kBeta,
                 thrust::raw_pointer_cast(means.data()), 1));
+        GpuErrorCheck(cudaDeviceSynchronize());
         return means;
     } else if (axis == 1) {
         thrust::device_vector<float> ones(M, 1.f);
@@ -67,6 +68,7 @@ thrust::device_vector<float> tsnecuda::utils::ReduceAlpha(
                 thrust::raw_pointer_cast(d_matrix.data()), N,
                 thrust::raw_pointer_cast(ones.data()), 1, &kBeta,
                 thrust::raw_pointer_cast(means.data()), 1));
+        GpuErrorCheck(cudaDeviceSynchronize());
         return means;
     } else {
         throw std::runtime_error("Axis must be 0 or 1.");

tsne/HIP/CMakeLists.txt

@@ -47,12 +47,25 @@ set(CMAKE_MODULE_PATH "${ROCM_PATH}/hip/cmake" ${CMAKE_MODULE_PATH})
 set(HIP_INCLUDE_DIRS  "${ROCM_PATH}/include"   ${HIP_INCLUDE_DIRS})
 set(HIP_LIBRARIES     "${ROCM_PATH}/lib"       ${HIP_LIBRARIES})
 
-# CMAKE_CXX_FLAGS
-if("${CMAKE_CXX_FLAGS}" STREQUAL "")
-    message(STATUS "Using DEFAULT compilation flags for the application")
-    set(CMAKE_CXX_FLAGS ${CMAKE_CXX_FLAGS} " -O3 -std=c++17 -ffast-math -D__HIP_PLATFORM_AMD__ ")
-else()
-    message(STATUS "OVERRIDING compilation flags")
+set(DEF_WL_CXX_FLAGS           " -D__HIP_PLATFORM_AMD__ ")
+set(DEF_GENERAL_CXX_FLAGS      " -O3 -std=c++17 -ffast-math ")
+set(DEF_COMBINED_CXX_FLAGS     "${DEF_GENERAL_CXX_FLAGS} ${DEF_WL_CXX_FLAGS}")
+
+# -DCMAKE_CXX_FLAGS=" -blah -blah " overrides the default flags (BOTH general and WL specific)
+# -DOVERRIDE_GENERAL_CXX_FLAGS=" -blah -blah " overrides the general flags only (and not the workload specific flags)
+# passing in both CMAKE_CXX_FLAGS and OVERRIDE_GENERAL_CXX_FLAGS is not allowed, in order to prevent ambiguity
+
+if(NOT "${CMAKE_CXX_FLAGS}" STREQUAL "" AND NOT "${OVERRIDE_GENERAL_CXX_FLAGS}" STREQUAL "")
+    message(FATAL_ERROR "Both  CMAKE_CXX_FLAGS and OVERRIDE_GENERAL_CXX_FLAGS cannot be passed in together")
+elseif("${CMAKE_CXX_FLAGS}" STREQUAL "" AND "${OVERRIDE_GENERAL_CXX_FLAGS}" STREQUAL "")
+    message(STATUS "Using DEFAULT compilation flags")
+    set(CMAKE_CXX_FLAGS "${DEF_COMBINED_CXX_FLAGS}")
+elseif(NOT "${OVERRIDE_GENERAL_CXX_FLAGS}" STREQUAL "")
+    message(STATUS "OVERRIDING GENERAL compilation flags")
+    set(CMAKE_CXX_FLAGS "${OVERRIDE_GENERAL_CXX_FLAGS}")
+    string(APPEND CMAKE_CXX_FLAGS ${DEF_WL_CXX_FLAGS})
+elseif(NOT "${CMAKE_CXX_FLAGS}" STREQUAL "")
+    message(STATUS "OVERRIDING GENERAL and WORKLOAD SPECIFIC compilation flags")
 endif()
 
 find_package(HIP REQUIRED)

tsne/HIP/src/kernels/nbodyfft.cpp

@@ -309,6 +309,7 @@ void DFT2D1gpu(float* din, thrust::complex<float>* dout, int num_rows, int num_c
     thrust::complex<float> sum, twiddle;
     angle = -2.0f * PI * fdividef((float)i, (float)num_cols);
     sum = 0.0f;
+#pragma unroll
     for (int k = 0; k < num_cols; ++k) {
         // sincosf(angle * k, &sinf, &cosf);
         // twiddle = thrust::complex<float>(cosf, sinf);
@@ -333,6 +334,7 @@ void DFT2D2gpu(thrust::complex<float>* din, thrust::complex<float>* dout, int nu
     thrust::complex<float> sum, twiddle;
     angle = -2.0f * PI * fdividef((float)i, (float)num_cols);
     sum = 0.0f;
+#pragma unroll
     for (int k = 0; k < num_cols; ++k) {
         // sincosf(angle * k, &sinf, &cosf);
         // twiddle = thrust::complex<float>(cosf, sinf);
@@ -357,6 +359,7 @@ void iDFT2D1gpu(thrust::complex<float>* din, thrust::complex<float>* dout, int n
     thrust::complex<float> sum, twiddle;
     angle = 2.0f * PI * fdividef((float)i, (float)num_cols);
     sum = 0.0f;
+#pragma unroll
     for (int k = 0; k < num_cols; ++k) {
         // sincosf(angle * k, &sinf, &cosf);
         // twiddle = thrust::complex<float>(cosf, sinf);
@@ -385,6 +388,7 @@ void iDFT2D2gpu(thrust::complex<float>* din, float* dout, int num_rows, int num_
     thrust::complex<float> twiddle;
     angle = 2.0f * PI * fdividef((float)i, (float)num_cols);
     sum = 0.0f;
+#pragma unroll
     for (int k = 0; k < num_cols; ++k) {
         // sincosf(angle * k, &sinf, &cosf);
         // twiddle = thrust::complex<float>(cosf, sinf);

tsne/HIP/src/kernels/perplexity_search.cpp

@@ -154,12 +154,12 @@ void tsnecuda::SearchPerplexity(
 
         // compute entropy of current row
         row_sum = tsnecuda::utils::ReduceSum(handle, pij, num_neighbors, num_points, 0);
-        GpuErrorCheck(hipDeviceSynchronize());
+        // GpuErrorCheck(hipDeviceSynchronize());
 
         // compute negative entropy
         thrust::transform(pij.begin(), pij.end(), entropy.begin(), tsnecuda::utils::FunctionalEntropy());
         neg_entropy = tsnecuda::utils::ReduceAlpha(handle, entropy, num_neighbors, num_points, -1.0f, 0);
-        GpuErrorCheck(hipDeviceSynchronize());
+        // GpuErrorCheck(hipDeviceSynchronize());
 
         // binary search for beta
         hipLaunchKernelGGL(PerplexitySearchKernel, NBLOCKS2, BLOCKSIZE2, 0, 0, 

tsne/HIP/src/utils/reduce_utils.cpp

@@ -57,6 +57,7 @@ thrust::device_vector<float> tsnecuda::utils::ReduceAlpha(
                 thrust::raw_pointer_cast(d_matrix.data()), N,
                 thrust::raw_pointer_cast(ones.data()), 1, &kBeta,
                 thrust::raw_pointer_cast(means.data()), 1));
+        GpuErrorCheck(hipDeviceSynchronize());
         return means;
     } else if (axis == 1) {
         thrust::device_vector<float> ones(M, 1.f);
@@ -67,6 +68,7 @@ thrust::device_vector<float> tsnecuda::utils::ReduceAlpha(
                 thrust::raw_pointer_cast(d_matrix.data()), N,
                 thrust::raw_pointer_cast(ones.data()), 1, &kBeta,
                 thrust::raw_pointer_cast(means.data()), 1));
+        GpuErrorCheck(hipDeviceSynchronize());
         return means;
     } else {
         throw std::runtime_error("Axis must be 0 or 1.");

tsne/README.md

@@ -2,22 +2,12 @@
 
 tsne implements [FIt-SNE algorithm](https://github.com/KlugerLab/FIt-SNE) for various GPU architectures (original CUDA source code is from [here](https://github.com/CannyLab/tsne-cuda)).
 
-## Cloning
-
-To clone
-
-```
-git clone https://github.com/oneapi-src/Velocity-Bench.git
-```
-
 ## Supported versions
 
 - CUDA: The original code was obtained from [here](https://github.com/CannyLab/tsne-cuda)
 - SYCL: The CUDA code was migrated using Intel DPCT, and then the resulting code was modified to remove the dpct headers.
 - HIP: Created from CUDA version using hipify-perl script.
 
-# Current Version:
-- Initial release of the workload
 
 # Build Instructions
 

tsne/SYCL/CMakeLists.txt

@@ -40,31 +40,48 @@ option(ENABLE_KERNEL_PROFILING "Build using kernel profiling" OFF)
 option(GPU_AOT                 "Build AOT for Intel GPU"      OFF)
 option(USE_NVIDIA_BACKEND      "Build for NVIDIA backend"     OFF)
 option(USE_AMDHIP_BACKEND      "Build for AMD HIP backend"    OFF)
+option(USE_SM                  "Build for specific SM"        OFF)
 
 if(ENABLE_KERNEL_PROFILING)
     message("-- Enabling kernel profiling")
     add_compile_options(-DENABLE_KERNEL_PROFILING)
 endif()
 
-set(INTEL_GPU_CXX_FLAGS  " -O2 -std=c++17 -fsycl -ffast-math -Wall -Wextra -Wno-unused-parameter -Wno-sign-compare -Wno-unknown-pragmas -Wno-unused-local-typedef ")
-set(NVIDIA_GPU_CXX_FLAGS " -O3 -std=c++17 -fsycl -ffast-math -Wall -Wextra -Wno-unused-parameter -Wno-sign-compare -Wno-unknown-pragmas -Wno-unused-local-typedef ")
-set(AMD_GPU_CXX_FLAGS    " -O3 -std=c++17 -fsycl -ffast-math -Wall -Wextra -Wno-unused-parameter -Wno-sign-compare -Wno-unknown-pragmas -Wno-unused-local-typedef ")
+set(DEF_INTEL_WL_CXX_FLAGS  " ")
+set(DEF_NVIDIA_WL_CXX_FLAGS " ")
+set(DEF_AMD_WL_CXX_FLAGS    " ")
 
-set(USE_DEFAULT_FLAGS ON)
-if("${CMAKE_CXX_FLAGS}" STREQUAL "")
+set(DEF_INTEL_GENERAL_CXX_FLAGS  " -O2 -std=c++17 -fsycl -ffast-math -Wall -Wextra -Wno-unused-parameter -Wno-sign-compare -Wno-unknown-pragmas -Wno-unused-local-typedef ")
+set(DEF_NVIDIA_GENERAL_CXX_FLAGS " -O3 -std=c++17 -fsycl -ffast-math -Wall -Wextra -Wno-unused-parameter -Wno-sign-compare -Wno-unknown-pragmas -Wno-unused-local-typedef ")
+set(DEF_AMD_GENERAL_CXX_FLAGS    " -O3 -std=c++17 -fsycl -ffast-math -Wall -Wextra -Wno-unused-parameter -Wno-sign-compare -Wno-unknown-pragmas -Wno-unused-local-typedef ")
+
+# -DCMAKE_CXX_FLAGS=" -blah -blah " overrides the default flags (BOTH general and WL specific)
+# -DOVERRIDE_GENERAL_CXX_FLAGS=" -blah -blah " overrides the general flags only (and not the workload specific flags)
+# passing in both CMAKE_CXX_FLAGS and OVERRIDE_GENERAL_CXX_FLAGS is not allowed, in order to prevent ambiguity
+
+if(NOT "${CMAKE_CXX_FLAGS}" STREQUAL "" AND NOT "${OVERRIDE_GENERAL_CXX_FLAGS}" STREQUAL "")
+    message(FATAL_ERROR "Both  CMAKE_CXX_FLAGS and OVERRIDE_GENERAL_CXX_FLAGS cannot be passed in together")
+elseif("${CMAKE_CXX_FLAGS}" STREQUAL "" AND "${OVERRIDE_GENERAL_CXX_FLAGS}" STREQUAL "")
     message(STATUS "Using DEFAULT compilation flags")
-else()
-    message(STATUS "OVERRIDING DEFAULT compilation flags")
-    set(USE_DEFAULT_FLAGS OFF)
+    set(INTEL_GPU_CXX_FLAGS        "${DEF_INTEL_GENERAL_CXX_FLAGS}   ${DEF_INTEL_WL_CXX_FLAGS}")
+    set(NVIDIA_GPU_CXX_FLAGS       "${DEF_NVIDIA_GENERAL_CXX_FLAGS}  ${DEF_NVIDIA_WL_CXX_FLAGS}")
+    set(AMD_GPU_CXX_FLAGS          "${DEF_AMD_GENERAL_CXX_FLAGS}     ${DEF_AMD_WL_CXX_FLAGS}")
+elseif(NOT "${OVERRIDE_GENERAL_CXX_FLAGS}" STREQUAL "")
+    message(STATUS "OVERRIDING GENERAL compilation flags")
+    set(INTEL_GPU_CXX_FLAGS        "${OVERRIDE_GENERAL_CXX_FLAGS}    ${DEF_INTEL_WL_CXX_FLAGS}")
+    set(NVIDIA_GPU_CXX_FLAGS       "${OVERRIDE_GENERAL_CXX_FLAGS}    ${DEF_NVIDIA_WL_CXX_FLAGS}")
+    set(AMD_GPU_CXX_FLAGS          "${OVERRIDE_GENERAL_CXX_FLAGS}    ${DEF_AMD_WL_CXX_FLAGS}")
+elseif(NOT "${CMAKE_CXX_FLAGS}" STREQUAL "")
+    message(STATUS "OVERRIDING GENERAL and WORKLOAD SPECIFIC compilation flags")
+    set(INTEL_GPU_CXX_FLAGS        "${CMAKE_CXX_FLAGS}")
+    set(NVIDIA_GPU_CXX_FLAGS       "${CMAKE_CXX_FLAGS}")
+    set(AMD_GPU_CXX_FLAGS          "${CMAKE_CXX_FLAGS}")
 endif()
 
-# JIT compilation 
 if(GPU_AOT)
     message(STATUS "Enabling INTEL backend")
-    if(USE_DEFAULT_FLAGS)
-        set(CMAKE_CXX_FLAGS "${INTEL_GPU_CXX_FLAGS}") # Default flags for Intel backend
-    endif()
-    if( (${GPU_AOT} STREQUAL "pvc") OR (${GPU_AOT} STREQUAL "PVC") )
+    set(CMAKE_CXX_FLAGS "${INTEL_GPU_CXX_FLAGS}")
+    if((${GPU_AOT} STREQUAL "pvc") OR (${GPU_AOT} STREQUAL "PVC"))
         message(STATUS "Enabling Intel GPU AOT compilation for ${GPU_AOT}")
         string(APPEND CMAKE_CXX_FLAGS " -fsycl-targets=spir64_gen -Xs \"-device 0x0bd5 -revision_id 0x2f\" ")
     else()
@@ -73,18 +90,28 @@ if(GPU_AOT)
     endif()
 elseif(USE_NVIDIA_BACKEND)
     message(STATUS "Enabling NVIDIA backend")
-    if(USE_DEFAULT_FLAGS)
-        set(CMAKE_CXX_FLAGS "${NVIDIA_GPU_CXX_FLAGS}") # Default flags for NV backend
+    set(CMAKE_CXX_FLAGS "${NVIDIA_GPU_CXX_FLAGS}")
+    if(USE_SM)
+        message("-- Building for SM_${USE_SM} compatibility")
+        string(APPEND CMAKE_CXX_FLAGS " -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend --cuda-gpu-arch=sm_${USE_SM} ")
+    else()
+        message("-- Building for SM_80 compatibility (DEFAULT)")
+        string(APPEND CMAKE_CXX_FLAGS " -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend --cuda-gpu-arch=sm_80 ")
     endif()
-    string(APPEND CMAKE_CXX_FLAGS " -fsycl-targets=nvptx64-nvidia-cuda ") # -O3 will be used, even though -O2 was set earlier 
+    string(APPEND CMAKE_CXX_FLAGS " -DUSE_NVIDIA_BACKEND")
 elseif(USE_AMDHIP_BACKEND)
     message(STATUS "Enabling AMD HIP backend for ${USE_AMDHIP_BACKEND} AMD architecture")
-    if(USE_DEFAULT_FLAGS)
-        set(CMAKE_CXX_FLAGS "${AMD_GPU_CXX_FLAGS}") # Default flags for AMD backend (gfx908 for MI100)
-    endif()
+    set(CMAKE_CXX_FLAGS "${AMD_GPU_CXX_FLAGS}")
     string(APPEND CMAKE_CXX_FLAGS " -fsycl-targets=amdgcn-amd-amdhsa -Xsycl-target-backend --offload-arch=${USE_AMDHIP_BACKEND} ")
+    string(APPEND CMAKE_CXX_FLAGS " -DUSE_AMDHIP_BACKEND")
+else()
+    # JIT case
+    message(STATUS "Enabling INTEL backend")
+    set(CMAKE_CXX_FLAGS "${INTEL_GPU_CXX_FLAGS}")
 endif()
 
+message(STATUS "CXX Compilation flags set to: ${CMAKE_CXX_FLAGS}")
+
 if(GPU_AOT)
     set(MKL_LINK static)
     set(MKL_THREADING sequential)