Merge branch 'master' into amd-master-next

Change-Id: Ib7e7824073f4dfc391fb3833fc90e11b327d3c22
2020-02-05 14:56:08 -08:00
@@ -30,7 +30,12 @@ THE SOFTWARE.
 __global__ void run_printf() { printf("Hello World\n"); }

 int main() {
-    hipLaunchKernelGGL(HIP_KERNEL_NAME(run_printf), dim3(1), dim3(1), 0, 0);
-    hipDeviceSynchronize();
+    int device_count = 0;
+    hipGetDeviceCount(&device_count);
+    for (int i = 0; i < device_count; ++i) {
+        hipSetDevice(i);
+        hipLaunchKernelGGL(HIP_KERNEL_NAME(run_printf), dim3(1), dim3(1), 0, 0);
+        hipDeviceSynchronize();
+    }
    passed();
 }
@@ -0,0 +1,212 @@
+/*
+Copyright (c) 2019 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 WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+// Simple test for hipLaunchCooperativeKernelMultiDevice API.
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM all
+ * TEST: %t
+ * HIT_END
+ */
+
+#include "hip/hip_runtime.h"
+#include "hip/hip_runtime_api.h"
+#include <iostream>
+#include <fstream>
+#include <vector>
+#include <thread>
+#include <chrono>
+#include "hip/hip_cooperative_groups.h"
+#include "test_common.h"
+
+using namespace std::chrono;
+
+const static uint NumOfLoopIterrations = 16 * 1024;
+const static uint BufferSizeInDwords = 28672 * NumOfLoopIterrations;
+const static uint numQueues = 4;
+const static uint numIter = 100;
+constexpr uint NumKernelArgs = 4;
+constexpr uint MaxGPUs = 8;
+
+#include <stdio.h>
+/*
+namespace cg = cooperative_groups;
+using namespace cooperative_groups;
+*/
+
+__global__ void test_gws(uint* buf, uint bufSize, long* tmpBuf, long* result)
+{
+    extern __shared__ long tmp[];
+    uint groups = gridDim.x;
+    uint group_id = blockIdx.x;
+    uint local_id = threadIdx.x;
+    uint chunk = gridDim.x * blockDim.x;
+
+    uint i = group_id * blockDim.x + local_id;
+    long sum = 0;
+    while (i < bufSize) {
+      sum += buf[i];
+      i += chunk;
+    }
+    tmp[local_id] = sum;
+    __syncthreads();
+    i = 0;
+    if (local_id == 0) {
+        sum = 0;
+        while (i < blockDim.x) {
+          sum += tmp[i];
+          i++;
+        }
+        tmpBuf[group_id] = sum;
+    }
+
+    // wait
+    cooperative_groups::this_grid().sync();
+
+    if (((blockIdx.x * blockDim.x) + threadIdx.x) == 0) {
+        for (uint i = 1; i < groups; ++i) {
+          sum += tmpBuf[i];
+       }
+       //*result = sum;
+       result[1 + cooperative_groups::this_multi_grid().grid_rank()] = sum;
+    }
+    cooperative_groups::this_multi_grid().sync();
+    if (cooperative_groups::this_multi_grid().grid_rank() == 0) {
+      sum = 0;
+      for (uint i = 1; i <= cooperative_groups::this_multi_grid().num_grids(); ++i) {
+        sum += result[i];
+      }
+      *result = sum;
+    }
+}
+
+int main() {
+  float *A, *B;
+  uint* dA[MaxGPUs];
+  long* dB[MaxGPUs];
+  long* dC;
+  hipModule_t Module;
+  hipStream_t stream[MaxGPUs];
+
+  uint32_t* init = new uint32_t[BufferSizeInDwords];
+  for (uint32_t i = 0; i < BufferSizeInDwords; ++i) {
+    init[i] = i;
+  }
+
+  int nGpu = 0;
+  HIPCHECK(hipGetDeviceCount(&nGpu));
+  size_t copySizeInDwords = BufferSizeInDwords / nGpu;
+  hipDeviceProp_t deviceProp[MaxGPUs];
+
+  for (int i = 0; i < nGpu; i++) {
+    HIPCHECK(hipSetDevice(i));
+
+    // Calculate the device occupancy to know how many blocks can be run concurrently
+    hipGetDeviceProperties(&deviceProp[i], 0);
+    if (!deviceProp[i].cooperativeMultiDeviceLaunch) {
+      printf("Device doesn't support cooperative launch!");
+      passed();
+      return 0;
+    }
+    size_t SIZE = copySizeInDwords * sizeof(uint);
+
+    HIPCHECK(hipMalloc((void**)&dA[i], SIZE));
+    if (i == 0) {
+      HIPCHECK(hipHostMalloc((void**)&dC, (nGpu + 1) * sizeof(long), hipHostMallocCoherent));
+    }
+    HIPCHECK(hipMemcpy(dA[i], &init[i * copySizeInDwords] , SIZE, hipMemcpyHostToDevice));
+    HIPCHECK(hipStreamCreate(&stream[i]));
+  }
+
+  dim3 dimBlock;
+  dim3 dimGrid;
+  dimGrid.x = 1;
+  dimGrid.y = 1;
+  dimGrid.z = 1;
+  dimBlock.x = 64;
+  dimBlock.y = 1;
+  dimBlock.z = 1;
+
+  int numBlocks = 0;
+  uint workgroups[3] = {64, 128, 256};
+
+  hipLaunchParams* launchParamsList = new hipLaunchParams[nGpu];
+
+  system_clock::time_point start = system_clock::now();
+
+  for (uint set = 0; set < 3; ++set) {
+    void* args[MaxGPUs * NumKernelArgs];
+    std::cout << "---------- Test#" << set << "---------------\n";
+    for (int i = 0; i < nGpu; i++) {
+      HIPCHECK(hipSetDevice(i));
+      dimBlock.x = workgroups[set];
+      HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks,
+      (hipFunction_t)test_gws, dimBlock.x * dimBlock.y * dimBlock.z, dimBlock.x * sizeof(long)));
+      
+      std::cout << "GPU(" << i << ") Block size: " << dimBlock.x << " Num blocks per CU: " << numBlocks << "\n";
+
+      dimGrid.x = deviceProp[i].multiProcessorCount * std::min(numBlocks, 32);
+      HIPCHECK(hipMalloc((void**)&dB[i], dimGrid.x * sizeof(long)));
+
+      args[i * NumKernelArgs]     = (void*)&dA[i];
+      args[i * NumKernelArgs + 1] = (void*)&copySizeInDwords;
+      args[i * NumKernelArgs + 2] = (void*)&dB[i];
+      args[i * NumKernelArgs + 3] = (void*)&dC;
+
+      launchParamsList[i].func = reinterpret_cast<void*>(test_gws);
+      launchParamsList[i].gridDim = dimGrid;
+      launchParamsList[i].blockDim = dimBlock;
+      launchParamsList[i].sharedMem = dimBlock.x * sizeof(long);
+      launchParamsList[i].stream = stream[i];
+      launchParamsList[i].args = &args[i * NumKernelArgs];
+    }
+ 
+    hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0);
+
+    HIPCHECK(hipMemcpy(init, dC, sizeof(long), hipMemcpyDeviceToHost));
+
+    if (*dC != (((long)(BufferSizeInDwords) * (BufferSizeInDwords - 1)) / 2)) {
+      std::cout << "Data validation failed for grid size = " << dimGrid.x << " and block size = " << dimBlock.x << "\n";
+      std::cout << "Test failed! \n";
+    }
+    for (int i = 0; i < nGpu; i++) {
+      hipFree(dB[i]);
+    }
+  }
+  system_clock::time_point end = system_clock::now();
+
+  delete [] launchParamsList;
+
+  std::chrono::duration<double> elapsed_seconds = end - start;
+
+  std::time_t end_time = std::chrono::system_clock::to_time_t(end);
+ 
+  std::cout << "finished computation at " << std::ctime(&end_time) <<
+    "elapsed time: " << elapsed_seconds.count() << "s\n";
+
+  hipSetDevice(0);
+  hipFree(dC);
+  for (int i = 0; i < nGpu; i++) {
+    hipFree(dA[i]);
+    HIPCHECK(hipStreamDestroy(stream[i]));
+  }
+  delete [] init;
+  passed();
+  return 0;
+}
@@ -116,7 +116,7 @@ int main() {

    dimBlock.x = workgroups[i];
    // Calculate the device occupancy to know how many blocks can be run concurrently
-    hipOccupancyMaxActiveBlocksPerMultiprocessor(reinterpret_cast<uint32_t*>(&numBlocks),
+    hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks,
        test_gws, dimBlock.x * dimBlock.y * dimBlock.z, dimBlock.x * sizeof(long));

    dimGrid.x = deviceProp.multiProcessorCount * std::min(numBlocks, 32);
@@ -30,10 +30,6 @@ THE SOFTWARE.
 #include "hip/hip_runtime.h"
 #include "test_common.h"

-#define fileName "vcpy_kernel.code"
-#define kernel_name "hello_world"
-
-
 __global__ void f1(float *a) { *a = 1.0; }

 template <typename T>
@@ -49,11 +45,10 @@ int main(int argc, char* argv[]) {
    hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, 0, 0);
    assert(gridSize != 0 && blockSize != 0);

-    uint32_t numBlock = 0;
-    hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, f1, blockSize, 0);
+    int numBlock = 0;
+    hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, f1, (int)blockSize, 0);
    assert(numBlock != 0);

-
    // test case for using kernel function pointer with template
    gridSize = 0;
    blockSize = 0;
@@ -61,17 +56,7 @@ int main(int argc, char* argv[]) {
    assert(gridSize != 0 && blockSize != 0);

    numBlock = 0;
-    hipOccupancyMaxActiveBlocksPerMultiprocessor<void(*)(int *)>(&numBlock, f2, blockSize, 0);
-    assert(numBlock != 0);
-
-
-    // test case for using kernel with hipFunction_t type
-    numBlock = 0;
-    hipModule_t Module;
-    hipFunction_t Function;
-    HIPCHECK(hipModuleLoad(&Module, fileName));
-    HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name));
-    HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, Function, blockSize, 0));
+    hipOccupancyMaxActiveBlocksPerMultiprocessor<void(*)(int *)>(&numBlock, f2, (int)blockSize, 0);
    assert(numBlock != 0);

    passed();
@@ -33,7 +33,6 @@ THE SOFTWARE.
 #define fileName "vcpy_kernel.code"
 #define kernel_name "hello_world"

-
 __global__ void f1(float *a) { *a = 1.0; }

 template <typename T>
@@ -76,8 +76,7 @@ void texture2Dtest()
    texDescr.readMode = hipReadModeElementType;

    hipTextureObject_t texObj;
-    hipResourceViewDesc resDesc;
-    HIPCHECK( hipCreateTextureObject(&texObj, &texRes, &texDescr, &resDesc));
+    HIPCHECK( hipCreateTextureObject(&texObj, &texRes, &texDescr, NULL));

    HIPCHECK(hipMalloc((void**)&devPtrB, SIZE_W*sizeof(TYPE_t)*SIZE_H)) ;