Merge 'master' into 'amd-master'

Change-Id: Ia02ee3c7c354a61f2c3027090ba2ee8d2803e37a


[ROCm/clr commit: 89c95bb227]

projects/clr/hipamd/bin/hipcc

@@ -320,7 +320,6 @@ my $runCmd = 1;
 my $buildDeps = 0;
 my $linkType = 1;
 my $setLinkType = 0;
-my $useCodeObjectV3 = 0;
 
 my @options = ();
 my @inputs  = ();
@@ -424,16 +423,6 @@ foreach $arg (@ARGV)
         $swallowArg = 1;
     }
 
-    if (($trimarg eq '-mno-code-object-v3') and ($HIP_PLATFORM eq 'clang') ) {
-        $useCodeObjectV3 = 0;
-        $swallowArg = 1;
-    }
-
-    if (($trimarg eq '-mcode-object-v3') and ($HIP_PLATFORM eq 'clang') ) {
-        $useCodeObjectV3 = 1;
-        $swallowArg = 1;
-    }
-
     if (($arg =~ /--genco/) and $HIP_PLATFORM eq 'clang' ) {
         $arg = "--cuda-device-only";
     }
@@ -767,12 +756,6 @@ if ($buildDeps and $HIP_PLATFORM eq 'clang') {
     $HIPCXXFLAGS .= " --cuda-host-only";
 }
 
-if ($useCodeObjectV3 and $HIP_PLATFORM eq 'clang') {
-    $HIPCXXFLAGS .= " -mcode-object-v3";
-} elsif ($HIP_PLATFORM eq 'clang') {
-    $HIPCXXFLAGS .= " -mno-code-object-v3";
-}
-
 # Add --hip-link only if there are no source files.
 if (!$needCXXFLAGS and $HIP_PLATFORM eq 'clang') {
     $HIPLDFLAGS .= " --hip-link";

projects/clr/hipamd/bin/hipconfig

@@ -59,6 +59,18 @@ sub parse_config_file {
     }
 }
 
+#---
+# Function to check if executable can be run
+sub can_run {
+    my ($exe) = @_;
+    `$exe --version 2>&1`;
+    if ($? == 0) {
+        return 1;
+    } else {
+        return 0;
+    }
+}
+
 $CUDA_PATH=$ENV{'CUDA_PATH'} // '/usr/local/cuda';
 $HCC_HOME=$ENV{'HCC_HOME'} // '/opt/rocm/hcc';
 $HSA_PATH=$ENV{'HSA_PATH'} // '/opt/rocm/hsa';
@@ -67,12 +79,13 @@ $HSA_PATH=$ENV{'HSA_PATH'} // '/opt/rocm/hsa';
 #HIP_PLATFORM controls whether to use NVCC or HCC for compilation:
 $HIP_PLATFORM=$ENV{'HIP_PLATFORM'};
 if (not defined $HIP_PLATFORM) {
-    $NAMDGPUNODES=`cat /sys/class/kfd/kfd/topology/nodes/*/properties 2>/dev/null | grep -c 'simd_count [1-9]'`;
+    if (can_run("$HCC_HOME/bin/hcc") or can_run("hcc")) {
-
+        $HIP_PLATFORM = "hcc";
-    if ($NAMDGPUNODES > 0) {
+    } elsif (can_run("$CUDA_PATH/bin/nvcc") or can_run("nvcc")) {
-        $HIP_PLATFORM = "hcc"
-    } else {
         $HIP_PLATFORM = "nvcc";
+    } else {
+        # Default to hcc for now
+        $HIP_PLATFORM = "hcc";
     }
 }
 

projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md

@@ -891,8 +891,8 @@
 | `cuMemsetD2D32Async`                                      |                               |
 | `cuMemsetD2D8`                                            |                               |
 | `cuMemsetD2D8Async`                                       |                               |
-| `cuMemsetD32`                                             | `hipMemset`                   |
+| `cuMemsetD32`                                             | `hipMemsetD32`                |
-| `cuMemsetD32Async`                                        | `hipMemsetAsync`              |
+| `cuMemsetD32Async`                                        | `hipMemsetD32Async`           |
 | `cuMemsetD8`                                              | `hipMemsetD8`                 |
 | `cuMemsetD8Async`                                         |                               |
 | `cuMipmappedArrayCreate`                                  |                               |

projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp

@@ -291,10 +291,10 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_FUNCTION_MAP{
   // no analogue
   {"cuMemsetD2D8Async",                                    {"hipMemsetD2D8Async",                                      "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}},
   // cudaMemset
-  {"cuMemsetD32",                                          {"hipMemset",                                               "", CONV_MEMORY, API_DRIVER}},
+  {"cuMemsetD32",                                          {"hipMemsetD32",                                            "", CONV_MEMORY, API_DRIVER}},
-  {"cuMemsetD32_v2",                                       {"hipMemset",                                               "", CONV_MEMORY, API_DRIVER}},
+  {"cuMemsetD32_v2",                                       {"hipMemsetD32",                                            "", CONV_MEMORY, API_DRIVER}},
   // cudaMemsetAsync
-  {"cuMemsetD32Async",                                     {"hipMemsetAsync",                                          "", CONV_MEMORY, API_DRIVER}},
+  {"cuMemsetD32Async",                                     {"hipMemsetD32Async",                                       "", CONV_MEMORY, API_DRIVER}},
   // no analogue
   {"cuMemsetD8",                                           {"hipMemsetD8",                                             "", CONV_MEMORY, API_DRIVER}},
   {"cuMemsetD8_v2",                                        {"hipMemsetD8",                                             "", CONV_MEMORY, API_DRIVER}},

projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h

@@ -1504,6 +1504,17 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes);
  */
 hipError_t hipMemsetD8(hipDeviceptr_t dest, unsigned char value, size_t sizeBytes);
 
+/**
+ *  @brief Fills the memory area pointed to by dest with the constant integer
+ * value for specified number of times.
+ *
+ *  @param[out] dst Data being filled
+ *  @param[in]  constant value to be set
+ *  @param[in]  number of values to be set
+ *  @return #hipSuccess, #hipErrorInvalidValue, #hipErrorNotInitialized
+ */
+hipError_t hipMemsetD32(hipDeviceptr_t dest, int value, size_t count);
+
 /**
  *  @brief Fills the first sizeBytes bytes of the memory area pointed to by dev with the constant
  * byte value value.
@@ -1521,6 +1532,24 @@ hipError_t hipMemsetD8(hipDeviceptr_t dest, unsigned char value, size_t sizeByte
  */
 hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream __dparm(0));
 
+/**
+ *  @brief Fills the memory area pointed to by dev with the constant integer
+ * value for specified number of times.
+ *
+ *  hipMemsetD32Async() is asynchronous with respect to the host, so the call may return before the
+ * memset is complete. The operation can optionally be associated to a stream by passing a non-zero
+ * stream argument. If stream is non-zero, the operation may overlap with operations in other
+ * streams.
+ *
+ *  @param[out] dst Pointer to device memory
+ *  @param[in]  value - Value to set for each byte of specified memory
+ *  @param[in]  count - number of values to be set
+ *  @param[in]  stream - Stream identifier
+ *  @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree
+ */
+hipError_t hipMemsetD32Async(hipDeviceptr_t dst, int value, size_t count,
+                             hipStream_t stream __dparm(0));
+
 /**
  *  @brief Fills the memory area pointed to by dst with the constant value.
  *

projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h

@@ -655,11 +655,20 @@ inline static hipError_t hipMemset(void* devPtr, int value, size_t count) {
     return hipCUDAErrorTohipError(cudaMemset(devPtr, value, count));
 }
 
+inline static hipError_t hipMemsetD32(hipDeviceptr_t devPtr, int value, size_t count) {
+    return hipCUResultTohipError(cuMemsetD32(devPtr, value, count));
+}
+
 inline static hipError_t hipMemsetAsync(void* devPtr, int value, size_t count,
                                         hipStream_t stream __dparm(0)) {
     return hipCUDAErrorTohipError(cudaMemsetAsync(devPtr, value, count, stream));
 }
 
+inline static hipError_t hipMemsetD32Async(hipDeviceptr_t devPtr, int value, size_t count,
+                                           hipStream_t stream __dparm(0)) {
+    return hipCUResultTohipError(cuMemsetD32Async(devPtr, value, count, stream));
+}
+
 inline static hipError_t hipMemsetD8(hipDeviceptr_t dest, unsigned char value, size_t sizeBytes) {
     return hipCUResultTohipError(cuMemsetD8(dest, value, sizeBytes));
 }

projects/clr/hipamd/src/hip_memory.cpp

@@ -350,7 +350,7 @@ hipError_t ihipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t heigh
      }
     // hardcoded 128 bytes
     *pitch = ((((int)width - 1) / 128) + 1) * 128;
-    const size_t sizeBytes = (*pitch) * height;
+    const size_t sizeBytes = (*pitch) * height * ((depth==0) ? 1 : depth);
 
     auto ctx = ihipGetTlsDefaultCtx();
 
@@ -1508,13 +1508,13 @@ __global__ void hip_copy2d_n(T* dst, const T* src, size_t width, size_t height,
 }  // namespace
 
 template <typename T>
-void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t sizeBytes) {
+void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t count) {
     static constexpr uint32_t block_dim = 256;
 
-    const uint32_t grid_dim = clamp_integer<size_t>(sizeBytes / block_dim, 1, UINT32_MAX);
+    const uint32_t grid_dim = clamp_integer<size_t>(count / block_dim, 1, UINT32_MAX);
 
     hipLaunchKernelGGL(hip_fill_n<block_dim>, dim3(grid_dim), dim3{block_dim}, 0u, stream, ptr,
-                       sizeBytes, std::move(val));
+                       count, std::move(val));
 }
 
 template <typename T>
@@ -1533,20 +1533,20 @@ typedef enum ihipMemsetDataType {
     ihipMemsetDataTypeInt    = 2
 }ihipMemsetDataType;
 
-hipError_t ihipMemset(void* dst, int  value, size_t sizeBytes, hipStream_t stream, enum ihipMemsetDataType copyDataType  )
+hipError_t ihipMemset(void* dst, int  value, size_t count, hipStream_t stream, enum ihipMemsetDataType copyDataType  )
 {
     hipError_t e = hipSuccess;
 
-    if (sizeBytes == 0) return e;
+    if (count == 0) return e;
 
     if (stream && (dst != NULL)) {
         if(copyDataType == ihipMemsetDataTypeChar){
-            if ((sizeBytes & 0x3) == 0) {
+            if ((count & 0x3) == 0) {
                 // use a faster dword-per-workitem copy:
                 try {
                     value = value & 0xff;
                     uint32_t value32 = (value << 24) | (value << 16) | (value << 8) | (value) ;
-                    ihipMemsetKernel<uint32_t> (stream, static_cast<uint32_t*> (dst), value32, sizeBytes/sizeof(uint32_t));
+                    ihipMemsetKernel<uint32_t> (stream, static_cast<uint32_t*> (dst), value32, count/sizeof(uint32_t));
                 }
                 catch (std::exception &ex) {
                     e = hipErrorInvalidValue;
@@ -1554,7 +1554,7 @@ hipError_t ihipMemset(void* dst, int  value, size_t sizeBytes, hipStream_t strea
              } else {
                 // use a slow byte-per-workitem copy:
                 try {
-                    ihipMemsetKernel<char> (stream, static_cast<char*> (dst), value, sizeBytes);
+                    ihipMemsetKernel<char> (stream, static_cast<char*> (dst), value, count);
                 }
                 catch (std::exception &ex) {
                     e = hipErrorInvalidValue;
@@ -1563,14 +1563,14 @@ hipError_t ihipMemset(void* dst, int  value, size_t sizeBytes, hipStream_t strea
         } else {
            if(copyDataType == ihipMemsetDataTypeInt) { // 4 Bytes value
                try {
-                   ihipMemsetKernel<uint32_t> (stream, static_cast<uint32_t*> (dst), value, sizeBytes);
+                   ihipMemsetKernel<uint32_t> (stream, static_cast<uint32_t*> (dst), value, count);
                } catch (std::exception &ex) {
                    e = hipErrorInvalidValue;
                }
             } else if(copyDataType == ihipMemsetDataTypeShort) {
                try {
                    value = value & 0xffff;
-                   ihipMemsetKernel<uint16_t> (stream, static_cast<uint16_t*> (dst), value, sizeBytes);
+                   ihipMemsetKernel<uint16_t> (stream, static_cast<uint16_t*> (dst), value, count);
                } catch (std::exception &ex) {
                    e = hipErrorInvalidValue;
                }
@@ -1719,6 +1719,18 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t st
     return ihipLogStatus(e);
 };
 
+hipError_t hipMemsetD32Async(hipDeviceptr_t dst, int value, size_t count, hipStream_t stream) {
+    HIP_INIT_SPECIAL_API(hipMemsetD32Async, (TRACE_MCMD), dst, value, count, stream);
+
+    hipError_t e = hipSuccess;
+
+    stream = ihipSyncAndResolveStream(stream);
+
+    e = ihipMemset(dst, value, count, stream, ihipMemsetDataTypeInt);
+
+    return ihipLogStatus(e);
+};
+
 hipError_t hipMemset(void* dst, int value, size_t sizeBytes) {
     HIP_INIT_SPECIAL_API(hipMemset, (TRACE_MCMD), dst, value, sizeBytes);
 
@@ -1787,6 +1799,22 @@ hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes
     return ihipLogStatus(e);
 }
 
+hipError_t hipMemsetD32(hipDeviceptr_t dst, int value, size_t count) {
+    HIP_INIT_SPECIAL_API(hipMemsetD32, (TRACE_MCMD), dst, value, count);
+
+    hipError_t e = hipSuccess;
+
+    hipStream_t stream = hipStreamNull;
+    stream = ihipSyncAndResolveStream(stream);
+    if (stream) {
+        e = ihipMemset(dst, value, count, stream, ihipMemsetDataTypeInt);
+        stream->locked_wait();
+    } else {
+        e = hipErrorInvalidValue;
+    }
+    return ihipLogStatus(e);
+}
+
 hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int  value, hipExtent extent )
 {
     HIP_INIT_SPECIAL_API(hipMemset3D, (TRACE_MCMD), &pitchedDevPtr, value, &extent);

projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset.cpp

@@ -26,43 +26,133 @@ THE SOFTWARE.
  * BUILD: %t %s ../../test_common.cpp
  * RUN: %t
  * //Small copy
- * RUN: %t -N 10    --memsetval 0x42
+ * RUN: %t -N 10    --memsetval 0x42 --memsetD32val 0x101
  * // Oddball size
- * RUN: %t -N 10013 --memsetval 0x5a
+ * RUN: %t -N 10013 --memsetval 0x5a --memsetD32val 0xDEADBEEF
  * // Big copy
- * RUN: %t -N 256M  --memsetval 0xa6
+ * RUN: %t -N 256M  --memsetval 0xa6 --memsetD32val 0xCAFEBABE
  * HIT_END
  */
 
 #include "hip/hip_runtime.h"
 #include "test_common.h"
 
+bool testhipMemset(int memsetval,int p_gpuDevice)
+{
+    size_t Nbytes = N*sizeof(char);
+    printf ("testhipMemset N=%zu  memsetval=%2x device=%d\n", N, memsetval, p_gpuDevice);
+    char *A_d;
+    char *A_h;
+    bool testResult = true;
 
-int main(int argc, char* argv[]) {
+    HIPCHECK ( hipMalloc(&A_d, Nbytes) );
-    HipTest::parseStandardArguments(argc, argv, true);
-
-    HIPCHECK(hipSetDevice(p_gpuDevice));
-
-    size_t Nbytes = N * sizeof(char);
-
-    printf("N=%zu  memsetval=%2x device=%d\n", N, memsetval, p_gpuDevice);
-
-    char* A_d;
-    char* A_h;
-
-    HIPCHECK(hipMalloc(&A_d, Nbytes));
     A_h = (char*)malloc(Nbytes);
+    HIPCHECK ( hipMemset(A_d, memsetval, Nbytes) );
+    HIPCHECK ( hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
 
-    HIPCHECK(hipMemset(A_d, memsetval, Nbytes));
+    for (int i=0; i<N; i++) {
-
-    HIPCHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
-
-    for (int i = 0; i < N; i++) {
         if (A_h[i] != memsetval) {
-            failed("mismatch at index:%d computed:%02x, memsetval:%02x\n", i, (int)A_h[i],
+            testResult = false;
-                   (int)memsetval);
+            printf("mismatch at index:%d computed:%02x, memsetval:%02x\n", i, (int)A_h[i], (int)memsetval);
+            break;
         }
     }
-
+    HIPCHECK(hipFree(A_d));
-    passed();
+    free(A_h);
+    return testResult;
+}
+
+bool testhipMemsetD32(int memsetD32val,int p_gpuDevice)
+{
+    size_t Nbytes = N*sizeof(int);
+    printf ("testhipMemsetD32 N=%zu  memsetD32val=%8x device=%d\n", N, memsetD32val, p_gpuDevice);
+    int *A_d;
+    int *A_h;
+    bool testResult = true;
+
+    HIPCHECK ( hipMalloc(&A_d, Nbytes) );
+    A_h = (int*)malloc(Nbytes);
+    HIPCHECK ( hipMemsetD32((hipDeviceptr_t)A_d, memsetD32val, N) );
+    HIPCHECK ( hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
+
+    for (int i=0; i<N; i++) {
+        if (A_h[i] != memsetD32val) {
+            testResult = false;            printf("mismatch at index:%d computed:%08x, memsetD32val:%08x\n", i, A_h[i], memsetD32val);
+            break;
+        }
+    }
+    HIPCHECK(hipFree(A_d));
+    free(A_h);
+    return testResult;
+}
+
+bool testhipMemsetAsync(int memsetval,int p_gpuDevice)
+{
+    size_t Nbytes = N*sizeof(int);
+    printf ("testhipMemsetAsync N=%zu  memsetval=%2x device=%d\n", N, memsetval, p_gpuDevice);
+    char *A_d;
+    char *A_h;
+    bool testResult = true;
+
+    HIPCHECK ( hipMalloc((void**)&A_d, Nbytes) );
+    A_h = (char*)malloc(Nbytes);
+    hipStream_t stream;
+    HIPCHECK(hipStreamCreate(&stream));
+    HIPCHECK ( hipMemsetAsync(A_d, memsetval, Nbytes, stream ));
+    HIPCHECK ( hipStreamSynchronize(stream));
+    HIPCHECK ( hipMemcpy(A_h, (void*)A_d, Nbytes, hipMemcpyDeviceToHost));
+
+    for (int i=0; i<N; i++) {
+        if (A_h[i] != memsetval) {
+            testResult = false;
+            printf("mismatch at index:%d computed:%02x, memsetval:%02x\n", i, (int)A_h[i], (int)memsetval);
+            break;
+        }
+    }
+    HIPCHECK(hipFree((void*)A_d));
+    HIPCHECK(hipStreamDestroy(stream));
+    free(A_h);
+    return testResult;
+}
+
+bool testhipMemsetD32Async(int memsetD32val,int p_gpuDevice)
+{
+    size_t Nbytes = N*sizeof(int);
+    printf ("testhipMemsetD32Async N=%zu  memsetval=%8x device=%d\n", N, memsetD32val, p_gpuDevice);
+    int *A_d;
+    int *A_h;
+    bool testResult = true;
+
+    HIPCHECK ( hipMalloc((void**)&A_d, Nbytes) );
+    A_h = (int*)malloc(Nbytes);
+    hipStream_t stream;
+    HIPCHECK(hipStreamCreate(&stream));
+    HIPCHECK ( hipMemsetD32Async((hipDeviceptr_t)A_d, memsetD32val, N, stream ));
+    HIPCHECK ( hipStreamSynchronize(stream));
+    HIPCHECK ( hipMemcpy(A_h, (void*)A_d, Nbytes, hipMemcpyDeviceToHost));
+
+    for (int i=0; i<N; i++) {
+        if (A_h[i] != memsetD32val) {
+            testResult = false;
+            printf("mismatch at index:%d computed:%02x, memsetD32val:%02x\n", i, A_h[i], memsetD32val);
+            break;
+        }
+    }
+    HIPCHECK(hipFree((void*)A_d));
+    HIPCHECK(hipStreamDestroy(stream));
+    free(A_h);
+    return testResult;
+}
+
+int main(int argc, char *argv[])
+{
+    HipTest::parseStandardArguments(argc, argv, true);
+    bool testResult = true;
+    HIPCHECK(hipSetDevice(p_gpuDevice));
+    testResult &= testhipMemset(memsetval, p_gpuDevice);
+    testResult &= testhipMemsetAsync(memsetval, p_gpuDevice);
+    testResult &= testhipMemsetD32(memsetD32val, p_gpuDevice);
+    testResult &= testhipMemsetD32Async(memsetD32val, p_gpuDevice);
+    if (testResult) passed();
+    failed("Output Mismatch\n"); 
 }

projects/clr/hipamd/tests/src/test_common.cpp

@@ -24,6 +24,7 @@ THE SOFTWARE.
 // standard global variables that can be set on command line
 size_t N = 4 * 1024 * 1024;
 char memsetval = 0x42;
+int memsetD32val = 0xDEADBEEF;
 int iterations = 1;
 unsigned blocksPerCU = 6;  // to hide latency
 unsigned threadsPerBlock = 256;
@@ -99,6 +100,12 @@ int parseStandardArguments(int argc, char* argv[], bool failOnUndefinedArg) {
                 failed("Bad memsetval argument");
             }
             memsetval = ex;
+        } else if (!strcmp(arg, "--memsetD32val")) {
+            int ex;
+            if (++i >= argc || !HipTest::parseInt(argv[i], &ex)) {
+                failed("Bad memsetD32val argument");
+            }
+            memsetD32val = ex;
         } else if (!strcmp(arg, "--iterations") || (!strcmp(arg, "-i"))) {
             if (++i >= argc || !HipTest::parseInt(argv[i], &iterations)) {
                 failed("Bad iterations argument");

projects/clr/hipamd/tests/src/test_common.h

@@ -98,6 +98,7 @@ THE SOFTWARE.
 // standard command-line variables:
 extern size_t N;
 extern char memsetval;
+extern int memsetD32val;
 extern int iterations;
 extern unsigned blocksPerCU;
 extern unsigned threadsPerBlock;