Merge 'master' into 'amd-master'

Change-Id: I2cda8cde846bc575585e988573ea70b03123ee4b [ROCm/hip commit: bb3b536834]
2019-05-24 07:27:12 -04:00
parent 8b82a2c5ec da977db4df
commit f16e39e322
15 changed files with 169 additions and 38 deletions
@@ -370,7 +370,6 @@ if($HIP_PLATFORM eq "nvcc"){

 my $toolArgs = "";  # arguments to pass to the hcc or nvcc tool
 my $optArg = ""; # -O args
-my $rdc = 0;

 foreach $arg (@ARGV)
 {
@@ -466,14 +465,6 @@ foreach $arg (@ARGV)
    {
        $optArg = $arg;
    }
-    if($arg =~ /-fgpu-rdc/)
-    {
-        $rdc = 1;
-    }
-    if($arg =~ /-fno-gpu-rdc/)
-    {
-        $rdc = 0;
-    }

    ## process linker response file for hip-clang
    ## extract object files from static library and pass them directly to
@@ -804,11 +795,7 @@ if ($HIP_PLATFORM eq "clang") {
        $HIPLDFLAGS .= " -O3";
    }
    $HIP_DEVLIB_FLAGS = " --hip-device-lib-path=$DEVICE_LIB_PATH";
-    if ($rdc eq 0) {
-        $HIPCXXFLAGS .= " $HIP_DEVLIB_FLAGS";
-    } else {
-        $HIPLDFLAGS .= " $HIP_DEVLIB_FLAGS";
-    }
+    $HIPCXXFLAGS .= " $HIP_DEVLIB_FLAGS";
    if ($isWindows) {
        $HIPCXXFLAGS .= " -std=c++14 -fms-extensions -fms-compatibility";
    } else {
@@ -848,7 +835,21 @@ if ($runCmd) {
        print ("HIP ($HIP_PATH) was built using hcc $hipConfig{'HCC_VERSION'}, but you are using $HCC_HOME/hcc with version $HCC_VERSION from hipcc. Please rebuild HIP including cmake or update HCC_HOME variable.\n") ;
        die unless $ENV{'HIP_IGNORE_HCC_VERSION'};
    }
-    system ("$CMD") or delete_temp_dirs () and die ();
+    system ("$CMD");
+    if ($? == -1) {
+        print "failed to execute: $!\n";
+        exit($?);
+    }
+    elsif ($? & 127) {
+        printf "child died with signal %d, %s coredump\n",
+        ($? & 127),  ($? & 128) ? 'with' : 'without';
+        exit($?);
+    }
+    else {
+         $CMD_EXIT_CODE = $? >> 8;
+    }
+    $? or delete_temp_dirs ();
+    exit($CMD_EXIT_CODE);
 }

 # vim: ts=4:sw=4:expandtab:smartindent
@@ -861,7 +861,7 @@
 | `cuMemAllocManaged`                                       |                               |
 | `cuMemAllocPitch`                                         |                               |
 | `cuMemcpy`                                                |                               |
-| `cuMemcpy2D`                                              |                               |
+| `cuMemcpy2D`                                              | `hipMemcpyParam2D`            |
 | `cuMemcpy2DAsync`                                         |                               |
 | `cuMemcpy2DUnaligned`                                     |                               |
 | `cuMemcpy3D`                                              |                               |
@@ -35,6 +35,9 @@
 - [Warp Cross-Lane Functions](#warp-cross-lane-functions)
  * [Warp Vote and Ballot Functions](#warp-vote-and-ballot-functions)
  * [Warp Shuffle Functions](#warp-shuffle-functions)
+- [Cooperative Groups Functions](#cooperative-groups-functions)
+- [Warp Matrix Functions](#warp-matrix-functions)
+- [Independent Thread Scheduling](#independent-thread-scheduling)
 - [Profiler Counter Function](#profiler-counter-function)
 - [Assert](#assert)
 - [Printf](#printf)
@@ -599,6 +602,70 @@ float __shfl_xor  (float var, int laneMask, int width=warpSize);

 ```

+## Cooperative Groups Functions
+
+Cooperative groups is a mechanism for forming and communicating between groups of threads at
+a granularity different than the block.  This feature was introduced in Cuda 9.
+
+HIP does not support any of the kernel language cooperative groups
+types or functions.
+
+
+| **Function** | **Supported in HIP** | **Supported in CUDA** |
+| --- | --- | --- |
+| `void thread_group.sync()` | | ✓ |
+| `unsigned thread_group.size()` | | ✓ |
+| `unsigned thread_group.thread_rank()` | | ✓ |
+| `bool thread_group.is_valid()` | | ✓ |
+| `thread_group tiled_partition(thread_group, size)` | | ✓ |
+| `thread_block_tile<N> tiled_partition<N>(thread_group)` | | ✓ |
+| `thread_block this_thread_block()` | | ✓ |
+| `T thread_block_tile.shfl()` | | ✓ |
+| `T thread_block_tile.shfl_down()` | | ✓ |
+| `T thread_block_tile.shfl_up()` | | ✓ |
+| `T thread_block_tile.shfl_xor()` | | ✓ |
+| `T thread_block_tile.any()` | | ✓ |
+| `T thread_block_tile.all()` | | ✓ |
+| `T thread_block_tile.ballot()` | | ✓ |
+| `T thread_block_tile.match_any()` | | ✓ |
+| `T thread_block_tile.match_all()` | | ✓ |
+| `coalesced_group coalesced_threads()` | | ✓ |
+| `grid_group this_grid()` | | ✓ |
+| `void grid_group.sync()` | | ✓ |
+| `unsigned grid_group.size()` | | ✓ |
+| `unsigned grid_group.thread_rank()` | | ✓ |
+| `bool grid_group.is_valid()` | | ✓ |
+| `multi_grid_group this_multi_grid()` | | ✓ |
+| `void multi_grid_group.sync()` | | ✓ |
+| `unsigned multi_grid_group.size()` | | ✓ |
+| `unsigned multi_grid_group.thread_rank()` | | ✓ |
+| `bool multi_grid_group.is_valid()` | | ✓ |
+
+## Warp Matrix Functions
+
+Warp matrix functions allow a warp to cooperatively operate on small matrices
+whose elements are spread over the lanes in an unspecified manner.  This feature
+was introduced in Cuda 9.
+
+HIP does not support any of the kernel language warp matrix
+types or functions.
+
+| **Function** | **Supported in HIP** | **Supported in CUDA** |
+| --- | --- | --- |
+| `void load_matrix_sync(fragment<...> &a, const T* mptr, unsigned lda)` | | ✓ |
+| `void load_matrix_sync(fragment<...> &a, const T* mptr, unsigned lda, layout_t layout)` | | ✓ |
+| `void store_matrix_sync(T* mptr, fragment<...> &a,  unsigned lda, layout_t layout)` | | ✓ |
+| `void fill_fragment(fragment<...> &a, const T &value)` | | ✓ |
+| `void mma_sync(fragment<...> &d, const fragment<...> &a, const fragment<...> &b, const fragment<...> &c , bool sat)` | | ✓ |
+
+## Independent Thread Scheduling
+
+The hardware support for independent thread scheduling introduced in certain architectures
+supporting Cuda allows threads to progress independently of each other and enables
+intra-warp synchronizations that were previously not allowed.
+
+HIP does not support this type of scheduling.
+
 ## Profiler Counter Function

 The Cuda `__prof_trigger()` instruction is not supported.
@@ -33,6 +33,7 @@ const std::map <llvm::StringRef, hipCounter> CUDA_INCLUDE_MAP{
  {"driver_types.h",            {"hip/driver_types.h",       "", CONV_INCLUDE,             API_RUNTIME}},
  {"cuda_fp16.h",               {"hip/hip_fp16.h",           "", CONV_INCLUDE,             API_RUNTIME}},
  {"cuda_texture_types.h",      {"hip/hip_texture_types.h",  "", CONV_INCLUDE,             API_RUNTIME}},
+  {"texture_fetch_functions.h", {"",                         "", CONV_INCLUDE,             API_RUNTIME}},
  {"vector_types.h",            {"hip/hip_vector_types.h",   "", CONV_INCLUDE,             API_RUNTIME}},
  {"cuda_profiler_api.h",       {"hip/hip_profile.h",        "", CONV_INCLUDE,             API_RUNTIME}},
  // cuComplex includes
@@ -175,8 +175,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_FUNCTION_MAP{
  {"cuMemcpy",                                             {"hipMemcpy_",                                              "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}},
  // no analogue
  // NOTE: Not equal to cudaMemcpy2D due to different signatures
-  {"cuMemcpy2D",                                           {"hipMemcpy2D_",                                            "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}},
-  {"cuMemcpy2D_v2",                                        {"hipMemcpy2D_",                                            "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}},
+  {"cuMemcpy2D",                                           {"hipMemcpyParam2D",                                        "", CONV_MEMORY, API_DRIVER}},
+  {"cuMemcpy2D_v2",                                        {"hipMemcpyParam2D",                                        "", CONV_MEMORY, API_DRIVER}},
  // no analogue
  // NOTE: Not equal to cudaMemcpy2DAsync due to different signatures
  {"cuMemcpy2DAsync",                                      {"hipMemcpy2DAsync_",                                       "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}},
@@ -189,6 +189,9 @@ bool HipifyAction::Exclude(const hipCounter & hipToken) {
      }
      return false;
    case CONV_INCLUDE:
+      if (hipToken.hipName.empty()) {
+        return true;
+      }
      switch (hipToken.apiType) {
        case API_RAND:
          if (hipToken.hipName == "hiprand_kernel.h") {
@@ -1009,14 +1009,20 @@ void __syncthreads()

 #define GETREG_IMMED(SZ,OFF,REG) (SZ << 11) | (OFF << 6) | REG

+/*
+  __smid returns the wave's assigned Compute Unit and Shader Engine.
+  The Compute Unit, CU_ID returned in bits 3:0, and Shader Engine, SE_ID in bits 5:4.
+  Note: the results vary over time.
+  SZ minus 1 since SIZE is 1-based.
+*/
 __device__
 inline
 unsigned __smid(void)
 {
    unsigned cu_id = __builtin_amdgcn_s_getreg(
-            GETREG_IMMED(HW_ID_CU_ID_SIZE, HW_ID_CU_ID_OFFSET, HW_ID));
+            GETREG_IMMED(HW_ID_CU_ID_SIZE-1, HW_ID_CU_ID_OFFSET, HW_ID));
    unsigned se_id = __builtin_amdgcn_s_getreg(
-            GETREG_IMMED(HW_ID_SE_ID_SIZE, HW_ID_SE_ID_OFFSET, HW_ID));
+            GETREG_IMMED(HW_ID_SE_ID_SIZE-1, HW_ID_SE_ID_OFFSET, HW_ID));

    /* Each shader engine has 16 CU */
    return (se_id << HW_ID_CU_ID_SIZE) + cu_id;
@@ -80,22 +80,22 @@ typedef struct hipArray {
 }hipArray;

 typedef struct hip_Memcpy2D {
-    size_t height;
-    size_t widthInBytes;
-    hipArray* dstArray;
-    hipDeviceptr_t dstDevice;
-    void* dstHost;
-    hipMemoryType dstMemoryType;
-    size_t dstPitch;
-    size_t dstXInBytes;
-    size_t dstY;
-    hipArray* srcArray;
-    hipDeviceptr_t srcDevice;
-    const void* srcHost;
-    hipMemoryType srcMemoryType;
-    size_t srcPitch;
    size_t srcXInBytes;
    size_t srcY;
+    hipMemoryType srcMemoryType;
+    const void* srcHost;
+    hipDeviceptr_t srcDevice;
+    hipArray* srcArray;
+    size_t srcPitch;
+    size_t dstXInBytes;
+    size_t dstY;
+    hipMemoryType dstMemoryType;
+    void* dstHost;
+    hipDeviceptr_t dstDevice;
+    hipArray* dstArray;
+    size_t dstPitch;
+    size_t WidthInBytes;
+    size_t Height;
 } hip_Memcpy2D;


@@ -1858,6 +1858,16 @@ hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc*
 */
 hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
                       size_t height, hipMemcpyKind kind);
+
+/**
+*  @brief Copies memory for 2D arrays.
+*  @param[in]   pCopy Parameters for the memory copy
+ *  @return      #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue,
+ * #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection
+ *
+ *  @see hipMemcpy, hipMemcpy2D, hipMemcpyToArray, hipMemcpy2DToArray, hipMemcpyFromArray,
+ * hipMemcpyToSymbol, hipMemcpyAsync
+*/
 hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy);

 /**
@@ -162,6 +162,7 @@ typedef CUdeviceptr hipDeviceptr_t;
 typedef struct cudaArray hipArray;
 typedef struct cudaArray* hipArray_const_t;
 typedef cudaFuncAttributes hipFuncAttributes;
+#define hip_Memcpy2D CUDA_MEMCPY2D
 #define hipMemcpy3DParms cudaMemcpy3DParms
 #define hipArrayDefault cudaArrayDefault
 #define hipArrayLayered cudaArrayLayered
@@ -578,6 +579,10 @@ inline static hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src,
        cudaMemcpy2D(dst, dpitch, src, spitch, width, height, hipMemcpyKindToCudaMemcpyKind(kind)));
 }

+inline static hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) {
+  return hipCUResultTohipError(cuMemcpy2D(pCopy));
+}
+
 inline static hipError_t hipMemcpy3D(const struct hipMemcpy3DParms *p)
 {
    return hipCUDAErrorTohipError(cudaMemcpy3D(p));
@@ -71,8 +71,8 @@ bool runTest(int argc, char** argv) {
    copyParam.srcMemoryType = hipMemoryTypeHost;
    copyParam.srcHost = hData;
    copyParam.srcPitch = width * sizeof(float);
-    copyParam.widthInBytes = copyParam.srcPitch;
-    copyParam.height = height;
+    copyParam.WidthInBytes = copyParam.srcPitch;
+    copyParam.Height = height;
    hipMemcpyParam2D(&copyParam);

    textureReference* texref;
@@ -1715,8 +1715,8 @@ hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) {
    if (pCopy == nullptr) {
        e = hipErrorInvalidValue;
    }
-    e = ihipMemcpy2D(pCopy->dstArray->data, pCopy->widthInBytes, pCopy->srcHost, pCopy->srcPitch,
-                     pCopy->widthInBytes, pCopy->height, hipMemcpyDefault);
+    e = ihipMemcpy2D(pCopy->dstArray->data, pCopy->WidthInBytes, pCopy->srcHost, pCopy->srcPitch,
+                     pCopy->WidthInBytes, pCopy->Height, hipMemcpyDefault);
    return ihipLogStatus(e);
 }

@@ -0,0 +1,36 @@
+// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
+/*
+Copyright (c) 2015-present Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+// CHECK: #include <hip/hip_runtime.h>
+#include <cuda.h>
+// CHECK-NOT: #include <texture_fetch_functions.h>
+#include <texture_fetch_functions.h>
+
+// CHECK: extern texture<float, 2, hipReadModeElementType> tex;
+extern texture<float, 2, cudaReadModeElementType> tex;
+
+extern "C" __global__ void tex2dKernel(float* outputData, int width, int height) {
+  int x = blockDim.x * blockIdx.x + threadIdx.x;
+  int y = blockDim.y * blockIdx.y + threadIdx.y;
+  outputData[y * width + x] = tex2D(tex, x, y);
+}
@@ -178,7 +178,9 @@ void runTests(int64_t numElements) {
    {
        test(0x01, C_d, C_h, numElements, syncNone, true /*expectMismatch*/);
        test(0x02, C_d, C_h, numElements, syncNullStream, false /*expectMismatch*/);
+#ifndef __HIP_CLANG_ONLY__
        test(0x04, C_d, C_h, numElements, syncOtherStream, true /*expectMismatch*/);
+#endif
        test(0x08, C_d, C_h, numElements, syncDevice, false /*expectMismatch*/);

        // Sending a marker to to null stream may synchronize the otherStream