Merge branch 'master' into tex_unbind_issue_fix

tests/hipify-clang/lit.cfg

@@ -84,17 +84,23 @@ clang_arguments = "-v"
 if sys.platform in ['win32']:
     run_test_ext = ".bat"
     hipify_path += "/" + config.build_type
+    # CUDA SDK ROOT
     clang_arguments += " -isystem'%s'/common/inc"
 else:
     run_test_ext = ".sh"
+    # CUDA SDK ROOT
     clang_arguments += " -isystem'%s'/samples/common/inc"
+# cuDNN ROOT
 clang_arguments += " -I'%s'/include"
 if config.pointer_size == 8:
     clang_arguments += " -D__LP64__"
 
+# CUB ROOT
+clang_arguments += " -I'%s'"
+
 hipify_arguments = "--cuda-path='%s'"
 
-config.substitutions.append(("%clang_args", clang_arguments % (config.cuda_sdk_root, config.cuda_dnn_root)))
+config.substitutions.append(("%clang_args", clang_arguments % (config.cuda_sdk_root, config.cuda_dnn_root, config.cuda_cub_root)))
 config.substitutions.append(("%hipify_args", hipify_arguments % (config.cuda_root)))
 config.substitutions.append(("hipify", '"' + hipify_path + "/hipify-clang" + '"'))
 config.substitutions.append(("%run_test", '"' + config.test_source_root + "/run_test" + run_test_ext + '"'))

tests/hipify-clang/lit.site.cfg.in

@@ -8,6 +8,7 @@ config.llvm_tools_dir = "@LLVM_TOOLS_BINARY_DIR@"
 config.obj_root = "@CMAKE_CURRENT_BINARY_DIR@"
 config.cuda_root = "@CUDA_TOOLKIT_ROOT_DIR@"
 config.cuda_dnn_root = "@CUDA_DNN_ROOT_DIR@"
+config.cuda_cub_root = "@CUDA_CUB_ROOT_DIR@"
 config.cuda_version_major = int("@CUDA_VERSION_MAJOR@")
 config.cuda_version_minor = int("@CUDA_VERSION_MINOR@")
 config.cuda_version = "@CUDA_VERSION@"

tests/hipify-clang/unit_tests/libraries/CUB/cub_01.cu

@@ -0,0 +1,60 @@
+// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
+// CHECK: #include <hip/hip_runtime.h>
+#include <iostream>
+// CHECK: #include <hiprand.h>
+#include <curand.h>
+// CHECK: #include <hipcub/hipcub.hpp>
+#include <cub/cub.cuh>
+
+#include <iostream>
+
+// TODO:
+// using namespace cub;
+
+template <typename T>
+__global__ void sort(const T* data_in, T* data_out){
+     // CHECK: typedef ::hipcub::BlockRadixSort<T, 1024, 4> BlockRadixSortT;
+     typedef ::cub::BlockRadixSort<T, 1024, 4> BlockRadixSortT;
+    __shared__ typename BlockRadixSortT::TempStorage tmp_sort;
+    double items[4];
+    int i0 = 4 * (blockIdx.x * blockDim.x + threadIdx.x);
+    for (int i = 0; i < 4; ++i){
+        items[i] = data_in[i0 + i];
+    }
+    BlockRadixSortT(tmp_sort).Sort(items);
+    for (int i = 0; i < 4; ++i){
+        data_out[i0 + i] = items[i];
+    }
+}
+
+int main(){
+    double* d_gpu = NULL;
+    double* result_gpu = NULL;
+    double* data_sorted = new double[4096];
+    // Allocate memory on the GPU
+    // CHECK: hipMalloc(&d_gpu, 4096 * sizeof(double));
+    cudaMalloc(&d_gpu, 4096 * sizeof(double));
+    // CHECK: hipMalloc(&result_gpu, 4096 * sizeof(double));
+    cudaMalloc(&result_gpu, 4096 * sizeof(double));
+    // CHECK: hiprandGenerator_t gen;
+    curandGenerator_t gen;
+    // Create generator
+    // CHECK: hiprandCreateGenerator(&gen, HIPRAND_RNG_PSEUDO_DEFAULT);
+    curandCreateGenerator(&gen, CURAND_RNG_PSEUDO_DEFAULT);
+    // Fill array with random numbers
+    // CHECK: hiprandGenerateNormalDouble(gen, d_gpu, 4096, 0.0, 1.0);
+    curandGenerateNormalDouble(gen, d_gpu, 4096, 0.0, 1.0);
+    // Destroy generator
+    // CHECK: hiprandDestroyGenerator(gen);
+    curandDestroyGenerator(gen);
+    // Sort data
+    // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(sort), dim3(1), dim3(1024), 0, 0, d_gpu, result_gpu);
+    sort<<<1, 1024>>>(d_gpu, result_gpu);
+    // CHECK: hipMemcpy(data_sorted, result_gpu, 4096 * sizeof(double), hipMemcpyDeviceToHost);
+    cudaMemcpy(data_sorted, result_gpu, 4096 * sizeof(double), cudaMemcpyDeviceToHost);
+    // Write the sorted data to standard out
+    for (int i = 0; i < 4096; ++i){
+        std::cout << data_sorted[i] << ", ";
+    }
+    std::cout << std::endl;
+}

tests/hipify-clang/unit_tests/samples/allocators.cu

@@ -1,9 +1,9 @@
 // RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
 
 #pragma once
-
+// CHECK: #include <hip/hip_runtime.h>
 #include <cuda_runtime.h>
-
+#include <math.h>
 
 /**
  * Allocate GPU memory for `count` elements of type `T`.
@@ -16,3 +16,38 @@ static T* gpuMalloc(size_t count) {
     return ret;
 }
 
+template<typename T>
+__global__ void add(int n, T* x, T* y) {
+  int index = blockIdx.x * blockDim.x + threadIdx.x;
+  int stride = blockDim.x * gridDim.x;
+  for (int i = index; i < n; i += stride)
+    y[i] = x[i] + y[i];
+}
+
+int main(int argc, char* argv[]) {
+    size_t numElements = 50;
+    float *A = gpuMalloc<float>(numElements);
+    float* B = gpuMalloc<float>(numElements);
+    for (int i = 0; i < numElements; ++i) {
+        A[i] = 1.0f;
+        B[i] = 2.0f;
+    }
+    int blockSize = 512;
+    int numBlocks = (numElements + blockSize - 1) / blockSize;
+    dim3 dimGrid(numBlocks, 1, 1);
+    dim3 dimBlock(blockSize, 1, 1);
+    // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(add<float>), dim3(dimGrid), dim3(dimBlock), 0, 0, numElements, A, B);
+    add<float><<<dimGrid, dimBlock>>>(numElements, A, B);
+    // CHECK: hipDeviceSynchronize();
+    cudaDeviceSynchronize();
+    float maxError = 0.0f;
+    for (int i = 0; i < numElements; ++i)
+      maxError = fmax(maxError, fabs(B[i] - 3.0f));
+    // CHECK: hipFree(A);
+    cudaFree(A);
+    // CHECK: hipFree(B);
+    cudaFree(B);
+    if (maxError == 0.0f)
+      return 0;
+    return -1;
+  }

tests/src/hiprtc/hiprtcGetLoweredName.cpp

@@ -20,7 +20,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 /* HIT_START
- * BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM all
+ * BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc
  * TEST: %t
  * HIT_END
  */
@@ -80,7 +80,16 @@ int main()
 
     for (auto&& x : variable_name_vec) hiprtcAddNameExpression(prog, x.c_str());
 
-    hiprtcResult compileResult = hiprtcCompileProgram(prog, 0, nullptr);
+    hipDeviceProp_t props;
+    int device = 0;
+    hipGetDeviceProperties(&props, device);
+    std::string gfxName = "gfx" + std::to_string(props.gcnArch);
+    std::string sarg = "--gpu-architecture=" + gfxName;
+    const char* options[] = {
+            sarg.c_str()
+    };
+
+    hiprtcResult compileResult = hiprtcCompileProgram(prog, 1, options);
 
     // Obtain compilation log from the program.
     size_t logSize;

tests/src/hiprtc/hiprtcGetTypeName.cpp

@@ -20,7 +20,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 /* HIT_START
- * BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM all
+ * BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc
  * TEST: %t
  * HIT_END
  */
@@ -76,7 +76,16 @@ int main()
 
     for (auto&& x : name_vec) hiprtcAddNameExpression(prog, x.c_str());
 
-    hiprtcResult compileResult = hiprtcCompileProgram(prog, 0, nullptr);
+    hipDeviceProp_t props;
+    int device = 0;
+    hipGetDeviceProperties(&props, device);
+    std::string gfxName = "gfx" + std::to_string(props.gcnArch);
+    std::string sarg = "--gpu-architecture=" + gfxName;
+    const char* options[] = {
+            sarg.c_str()
+    };
+
+    hiprtcResult compileResult = hiprtcCompileProgram(prog, 1, options);
 
     size_t logSize;
     hiprtcGetProgramLogSize(prog, &logSize);

tests/src/hiprtc/saxpy.cpp

@@ -20,7 +20,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 /* HIT_START
- * BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM all
+ * BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc
  * TEST: %t
  * HIT_END
  */
@@ -66,7 +66,16 @@ int main()
                         nullptr,    // headers
                         nullptr);   // includeNames
 
-    hiprtcResult compileResult{hiprtcCompileProgram(prog, 0, nullptr)};
+    hipDeviceProp_t props;
+    int device = 0;
+    hipGetDeviceProperties(&props, device);
+    std::string gfxName = "gfx" + std::to_string(props.gcnArch);
+    std::string sarg = "--gpu-architecture=" + gfxName;
+    const char* options[] = {
+        sarg.c_str()
+    };
+
+    hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, options)};
 
     size_t logSize;
     hiprtcGetProgramLogSize(prog, &logSize);

tests/src/runtimeApi/memory/hipArray.cpp

@@ -21,7 +21,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM all
  * TEST: %t
  * HIT_END
  */

tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp

@@ -22,7 +22,7 @@ THE SOFTWARE.
 // Simple test for hipLaunchCooperativeKernel API.
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM all
  * TEST: %t
  * HIT_END
  */