Merge 'amd-master-next' into 'amd-npi-next'

Change-Id: Idb00a49722b2df682d15451bcb6b1c6698e0a896

hipamd/CMakeLists.txt

@@ -2,7 +2,7 @@ cmake_minimum_required(VERSION 3.4.3)
 project(hip)
 # sample command for hip-rocclr, you'll need to have rocclr installed
 #  cmake -DHIP_COMPILER=clang -DHIP_PLATFORM=rocclr ..
-#  cmake -DHIP_COMPILER=clang -DHIP_PLATFORM=rocclr -DROCclr_DIR=/extra/lmoriche/hip-rocclr/rocclr -DOPENCL_DIR=/extra/lmoriche/clients/lmoriche_opencl_dev2/drivers/opencl/api/opencl -DLIBROCclr_STATIC_DIR=/extra/lmoriche/hip-rocclr/build/rocclr ..
+#  cmake -DHIP_COMPILER=clang -DHIP_PLATFORM=rocclr -DOPENCL_DIR=/path/to/opencl/api/opencl -DCMAKE_PREFIX_PATH=/path/to/rocclr/build/or/install/directory ..
 
 set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} "${CMAKE_CURRENT_SOURCE_DIR}/cmake")
 

hipamd/include/hip/hcc_detail/hip_cooperative_groups.h

@@ -166,6 +166,55 @@ this_grid() {
   return grid_group(internal::grid::size());
 }
 
+/** \brief The workgroup (thread-block in CUDA terminology) cooperative group
+ *         type
+ *
+ *  \details Represents an intra-workgroup cooperative group type where the
+ *           participating threads within the group are exctly the same threads
+ *           which are participated in the currently executing `workgroup`
+ */
+class thread_block : public thread_group {
+  // Only these friend functions are allowed to construct an object of this
+  // class and access its resources
+  friend __CG_QUALIFIER__ thread_block this_thread_block();
+
+ protected:
+  // Construct a workgroup thread group (through the API this_thread_block())
+  explicit __CG_QUALIFIER__ thread_block(uint32_t size)
+      : thread_group(internal::cg_workgroup, size) { }
+
+ public:
+  // 3-dimensional block index within the grid
+  __CG_QUALIFIER__ dim3 group_index() {
+    return internal::workgroup::group_index();
+  }
+  // 3-dimensional thread index within the block
+  __CG_QUALIFIER__ dim3 thread_index() {
+    return internal::workgroup::thread_index();
+  }
+  __CG_QUALIFIER__ uint32_t thread_rank() const {
+    return internal::workgroup::thread_rank();
+  }
+  __CG_QUALIFIER__ bool is_valid() const {
+    return internal::workgroup::is_valid();
+  }
+  __CG_QUALIFIER__ void sync() const {
+    internal::workgroup::sync();
+  }
+};
+
+/** \brief User exposed API interface to construct workgroup cooperative
+ *         group type object - `thread_block`
+ *
+ *  \details User is not allowed to directly construct an object of type
+ *           `thread_block`. Instead, he should construct it through this API
+ *           function
+ */
+__CG_QUALIFIER__ thread_block
+this_thread_block() {
+  return thread_block(internal::workgroup::size());
+}
+
 /**
  *  Implemenation of all publicly exposed base class APIs
  */
@@ -177,6 +226,9 @@ __CG_QUALIFIER__ uint32_t thread_group::thread_rank() const {
     case internal::cg_grid: {
       return (static_cast<const grid_group*>(this)->thread_rank());
     }
+    case internal::cg_workgroup: {
+      return (static_cast<const thread_block*>(this)->thread_rank());
+    }
     default: {
       return 0; //TODO(mahesha)
     }
@@ -191,6 +243,9 @@ __CG_QUALIFIER__ bool thread_group::is_valid() const {
     case internal::cg_grid: {
       return (static_cast<const grid_group*>(this)->is_valid());
     }
+    case internal::cg_workgroup: {
+      return (static_cast<const thread_block*>(this)->is_valid());
+    }
     default: {
       return false;
     }
@@ -207,6 +262,10 @@ __CG_QUALIFIER__ void thread_group::sync() const {
       static_cast<const grid_group*>(this)->sync();
       break;
     }
+    case internal::cg_workgroup: {
+      static_cast<const thread_block*>(this)->sync();
+      break;
+    }
   }
 }
 

hipamd/include/hip/hcc_detail/hip_cooperative_groups_helper.h

@@ -60,7 +60,8 @@ namespace internal {
 typedef enum {
   cg_invalid,
   cg_multi_grid,
-  cg_grid
+  cg_grid,
+  cg_workgroup
 } group_type;
 
 /**
@@ -136,6 +137,43 @@ __CG_STATIC_QUALIFIER__ void sync() {
 
 } // namespace grid
 
+/**
+ *  Functionalities related to `workgroup` (thread_block in CUDA terminology)
+ *  cooperative group type
+ */
+namespace workgroup {
+
+__CG_STATIC_QUALIFIER__ dim3 group_index() {
+  return (dim3((uint32_t)hipBlockIdx_x, (uint32_t)hipBlockIdx_y,
+               (uint32_t)hipBlockIdx_z));
+}
+
+__CG_STATIC_QUALIFIER__ dim3 thread_index() {
+  return (dim3((uint32_t)hipThreadIdx_x, (uint32_t)hipThreadIdx_y,
+               (uint32_t)hipThreadIdx_z));
+}
+
+__CG_STATIC_QUALIFIER__ uint32_t size() {
+  return((uint32_t)(hipBlockDim_x * hipBlockDim_y * hipBlockDim_z));
+}
+
+__CG_STATIC_QUALIFIER__ uint32_t thread_rank() {
+ return ((uint32_t)((hipThreadIdx_z * hipBlockDim_y * hipBlockDim_x) +
+                    (hipThreadIdx_y * hipBlockDim_x) +
+                    (hipThreadIdx_x)));
+}
+
+__CG_STATIC_QUALIFIER__ bool is_valid() {
+   //TODO(mahesha) any functionality need to be added here? I believe not
+  return true;
+}
+
+__CG_STATIC_QUALIFIER__ void sync() {
+  __syncthreads();
+}
+
+} // namespace workgroup
+
 } // namespace internal
 
 } // namespace cooperative_groups

hipamd/include/hip/hcc_detail/hip_runtime_api.h

@@ -849,6 +849,28 @@ hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int* flags);
 hipError_t hipStreamGetPriority(hipStream_t stream, int* priority);
 
 
+/**
+ * @brief Create an asynchronous stream with the specified CU mask.
+ *
+ * @param[in, out] stream Pointer to new stream
+ * @param[in ] cuMaskSize Size of CU mask bit array passed in.
+ * @param[in ] cuMask Bit-vector representing the CU mask. Each active bit represents using one CU.
+ * The first 32 bits represent the first 32 CUs, and so on. If its size is greater than physical
+ * CU number (i.e., multiProcessorCount member of hipDeviceProp_t), the extra elements are ignored.
+ * It is user's responsibility to make sure the input is meaningful.
+ * @return #hipSuccess, #hipErrorInvalidHandle, #hipErrorInvalidValue
+ *
+ * Create a new asynchronous stream with the specified CU mask.  @p stream returns an opaque handle
+ * that can be used to reference the newly created stream in subsequent hipStream* commands.  The
+ * stream is allocated on the heap and will remain allocated even if the handle goes out-of-scope.
+ * To release the memory used by the stream, application must call hipStreamDestroy.
+ *
+ *
+ * @see hipStreamCreate, hipStreamSynchronize, hipStreamWaitEvent, hipStreamDestroy
+ */
+hipError_t hipExtStreamCreateWithCUMask(hipStream_t* stream, uint32_t cuMaskSize, const uint32_t* cuMask);
+
+
 /**
  * Stream CallBack struct
  */

hipamd/rocclr/CMakeLists.txt

@@ -27,38 +27,11 @@ endif()
 
 set(USE_PROF_API "1")
 
-# FIXME: Make this required and remove the legacy handling below
-set(save_rocclr_dir ${ROCclr_DIR})
-set(save_rocclr_static_dir ${LIBROCclr_STATIC_DIR})
-
-find_package(ROCclr CONFIG
+find_package(ROCclr REQUIRED CONFIG
   PATHS
     /opt/rocm
     /opt/rocm/rocclr)
 
-if (NOT ROCclr_FOUND)
-  if(NOT DEFINED LIBROCclr_STATIC_DIR)
-    find_path(LIBROCclr_STATIC_DIR
-      NAMES libamdrocclr_static.a
-      PATHS /opt/rocm/rocclr
-      PATH_SUFFIXES lib)
-  else()
-    set(LIBROCclr_STATIC_DIR ${save_rocclr_static_dir})
-  endif()
-
-  if(NOT DEFINED ROCclr_DIR)
-    find_path(ROCclr_DIR
-      NAMES top.hpp
-      PATH_SUFFIXES include
-      PATHS /opt/rocm/rocclr)
-  else()
-    set(ROCclr_DIR ${save_rocclr_dir})
-  endif()
-  message("Found Static rocclr lib:${LIBROCclr_STATIC_DIR} and rocclr includes: ${ROCclr_DIR}")
-  include(${LIBROCclr_STATIC_DIR}/amdrocclr_staticTargets.cmake)
-endif()
-
-set(PROF_API_HEADER_PATH ${ROCclr_DIR}/platform)
 #############################
 # Profiling API support
 #############################
@@ -132,14 +105,6 @@ target_include_directories(hip64
     ${PROJECT_SOURCE_DIR}
     ${PROJECT_SOURCE_DIR}/amdocl
     ${PROJECT_SOURCE_DIR}/include/hip/hcc_detail/elfio
-
-    # FIXME: Remove ROCclr_DIr explicit references
-    ${ROCclr_DIR}
-    ${ROCclr_DIR}/include
-    ${ROCclr_DIR}/compiler/lib
-    ${ROCclr_DIR}/compiler/lib/include
-    ${ROCclr_DIR}/elf/utils/common
-    ${ROCclr_DIR}/elf/utils/libelf
     ${ROCR_INCLUDES}
     $<TARGET_PROPERTY:amdrocclr_static,INTERFACE_INCLUDE_DIRECTORIES>
     $<TARGET_PROPERTY:amd_comgr,INTERFACE_INCLUDE_DIRECTORIES>)
@@ -204,17 +169,6 @@ set_target_properties(hip64 PROPERTIES PUBLIC_HEADER ${PROF_API_STR})
 set_target_properties(amdhip64 PROPERTIES PUBLIC_HEADER ${PROF_API_STR})
 set_target_properties(amdhip64_static PROPERTIES PUBLIC_HEADER ${PROF_API_STR})
 
-
-
-# We expect amdhip64_static to contain objects of rocclr and hip. But linker
-# let amdhip64_static contain objects of hip only. So we will use a
-# a custom amdhip64_static_combiner to combine objects of vid and hip into
-# amdhip64_static. To avoid amdhip64_static contains itself,
-# amdhip64_static_temp is created internally.
-add_library(amdhip64_static_temp STATIC
-    $<TARGET_OBJECTS:hip64>
-    )
-
 add_library(host INTERFACE)
 target_link_libraries(host INTERFACE amdhip64)
 add_library(device INTERFACE)
@@ -226,19 +180,19 @@ target_link_libraries(device INTERFACE host)
 # filename.
 target_link_libraries(amdhip64 PRIVATE amdrocclr_static Threads::Threads dl)
 target_link_libraries(amdhip64_static PRIVATE Threads::Threads dl)
-target_link_libraries(amdhip64_static_temp PRIVATE Threads::Threads dl)
 
 # combine objects of vid and hip into amdhip64_static
 add_custom_target(
     amdhip64_static_combiner
     ALL
-    COMMAND rm -f $<TARGET_FILE:amdhip64_static>  # Must remove old one, otherwise the new one will contain obsolete stuff
-    COMMAND ${CMAKE_AR} -rcsT $<TARGET_FILE:amdhip64_static> $<TARGET_FILE:amdhip64_static_temp> $<TARGET_FILE:amdrocclr_static>
-    DEPENDS amdhip64_static amdhip64_static_temp amdrocclr_static # To make sure this is the last step
+    COMMAND rm -rf  static_lib_temp && mkdir static_lib_temp && cd static_lib_temp # Create temp folder to contain *.o
+    COMMAND ${CMAKE_AR} -x $<TARGET_FILE:amdrocclr_static> # Extract *.o from amdrocclr_static
+    COMMAND ${CMAKE_AR} -rcs $<TARGET_FILE:amdhip64_static> *.o # Append *.o to amdhip64_static
+    COMMAND cd .. && rm -rf  static_lib_temp # Remove temp folder
+    DEPENDS amdhip64_static amdrocclr_static # To make sure this is the last step
     COMMENT "Combining static libs into amdhip64_static"
 )
 
-
 INSTALL(PROGRAMS $<TARGET_FILE:amdhip64_static> DESTINATION lib COMPONENT MAIN)
 INSTALL(PROGRAMS $<TARGET_FILE:amdhip64> DESTINATION lib COMPONENT MAIN)
 INSTALL(CODE "execute_process( COMMAND ${CMAKE_COMMAND} -E create_symlink libamdhip64.so lib/libhip_hcc.so  )" DESTINATION lib COMPONENT MAIN)

hipamd/rocclr/hip_hcc.def.in

@@ -252,3 +252,4 @@ hipTexObjectDestroy
 hipTexObjectGetResourceDesc
 hipTexObjectGetResourceViewDesc
 hipTexObjectGetTextureDesc
+hipExtStreamCreateWithCUMask

hipamd/rocclr/hip_hcc.map.in

@@ -258,6 +258,7 @@ global:
     hipInitActivityCallback*;
     hipEnableActivityCallback*;
     hipGetCmdName*;
+    hipExtStreamCreateWithCUMask;
     };
 local:
     *;

hipamd/rocclr/hip_internal.hpp

@@ -87,9 +87,11 @@ namespace hip {
     amd::CommandQueue::Priority priority_;
     unsigned int flags_;
     bool null_;
+    const std::vector<uint32_t> cuMask_;
 
   public:
-    Stream(Device* dev, amd::CommandQueue::Priority p, unsigned int f = 0, bool null_stream = false);
+    Stream(Device* dev, amd::CommandQueue::Priority p, unsigned int f = 0, bool null_stream = false,
+           const std::vector<uint32_t>& cuMask = {});
 
     /// Creates the hip stream object, including AMD host queue
     bool Create();

hipamd/rocclr/hip_memory.cpp

@@ -1962,7 +1962,10 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* dev_ptr) {
 
   /* Create an handle for IPC. Store the memory size inside the handle */
   ihandle = reinterpret_cast<ihipIpcMemHandle_t *>(handle);
-  dev_mem_obj->IpcCreate(offset, &(ihandle->psize), &(ihandle->ipc_handle));
+  if(!dev_mem_obj->IpcCreate(offset, &(ihandle->psize), &(ihandle->ipc_handle))) {
+    DevLogPrintfError("IPC memory creation failed for memory: 0x%x", dev_ptr);
+    HIP_RETURN(hipErrorInvalidValue);
+  }
 
   HIP_RETURN(hipSuccess);
 }
@@ -2023,7 +2026,9 @@ hipError_t hipIpcCloseMemHandle(void* dev_ptr) {
   amd::MemObjMap::RemoveMemObj(amd_mem_obj);
 
   /* detach the memory */
-  device->IpcDetach(*amd_mem_obj);
+  if (!device->IpcDetach(*amd_mem_obj)){
+     HIP_RETURN(hipErrorInvalidHandle);
+  }
 
   HIP_RETURN(hipSuccess);
 }

hipamd/rocclr/hip_prof_gen.py

@@ -452,7 +452,12 @@ def generate_prof_header(f, api_map, opts_map):
         arg_type = arg_tuple[0]
         fld_name = arg_tuple[1]
         arg_name = opts_list[ind]
-        f.write('  cb_data.args.' + name + '.' + fld_name + ' = (' + arg_type + ')' + arg_name + '; \\\n')
+        if arg_type == "char*":
+          f.write('  cb_data.args.' + name + '.' + fld_name + ' = strdup("***"); \\\n')
+        elif arg_type == "const char*":
+          f.write('  cb_data.args.' + name + '.' + fld_name + ' = (' + arg_name + ') ? strdup(' + arg_name + ') : NULL; \\\n')
+        else:
+          f.write('  cb_data.args.' + name + '.' + fld_name + ' = (' + arg_type + ')' + arg_name + '; \\\n')
     f.write('};\n')
   f.write('#define INIT_CB_ARGS_DATA(cb_id, cb_data) INIT_##cb_id##_CB_ARGS_DATA(cb_data)\n')
 
@@ -467,13 +472,20 @@ def generate_prof_header(f, api_map, opts_map):
   f.write('  switch (id) {\n')
   for name, args in api_map.items():
     f.write('    case HIP_API_ID_' + name + ':\n')
-    f.write('      oss << "' + name + '("')
+    f.write('      oss << "' + name + '(";\n')
     for ind in range(0, len(args)):
       arg_tuple = args[ind]
+      arg_type = arg_tuple[0]
       arg_name = arg_tuple[1]
-      if ind != 0: f.write(' << ","')
-      f.write('\n          << " ' + arg_name  + '=" << data->args.' + name + '.' + arg_name)
-    f.write('\n          << ")";\n')
+      var_name = 'data->args.' + name + '.' + arg_name
+      delim = '' if ind == 0 else ', ';
+      oss_stream = 'oss << "' + delim + arg_name  + '='
+      line_shift = '      '
+      f.write(line_shift)
+      if re.search(r'\*', arg_type):
+        f.write('if (' + var_name + ' == NULL) ' + oss_stream + 'NULL";\n' + line_shift + 'else ')
+      f.write(oss_stream + '" << ' + var_name + ';\n')
+    f.write('      oss << ")";\n')
     f.write('    break;\n')
   f.write('    default: oss << "unknown";\n')
   f.write('  };\n')

hipamd/rocclr/hip_stream.cpp

@@ -44,15 +44,15 @@ namespace hip {
 
 // ================================================================================================
 Stream::Stream(hip::Device* dev, amd::CommandQueue::Priority p,
-    unsigned int f, bool null_stream)
+    unsigned int f, bool null_stream, const std::vector<uint32_t>& cuMask)
   : queue_(nullptr), lock_("Stream Callback lock"), device_(dev),
-    priority_(p), flags_(f), null_(null_stream) {}
+    priority_(p), flags_(f), null_(null_stream), cuMask_(cuMask) {}
 
 // ================================================================================================
 bool Stream::Create() {
   cl_command_queue_properties properties = CL_QUEUE_PROFILING_ENABLE;
   queue_ = new amd::HostQueue(*device_->asContext(), *device_->devices()[0], properties,
-                             amd::CommandQueue::RealTimeDisabled, priority_);
+                             amd::CommandQueue::RealTimeDisabled, priority_, cuMask_);
   // Create a host queue
   bool result = (queue_ != nullptr) ? queue_->create() : false;
   // Insert just created stream into the list of the blocking queues
@@ -160,8 +160,9 @@ void CL_CALLBACK ihipStreamCallback(cl_event event, cl_int command_exec_status,
 
 // ================================================================================================
 static hipError_t ihipStreamCreate(hipStream_t* stream,
-                                  unsigned int flags, amd::CommandQueue::Priority priority) {
-  hip::Stream* hStream = new hip::Stream(hip::getCurrentDevice(), priority, flags);
+                                  unsigned int flags, amd::CommandQueue::Priority priority,
+                                  const std::vector<uint32_t>& cuMask = {}) {
+  hip::Stream* hStream = new hip::Stream(hip::getCurrentDevice(), priority, flags, false, cuMask);
 
   if (hStream == nullptr) {
     return hipErrorOutOfMemory;
@@ -311,3 +312,21 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback
 
   HIP_RETURN(hipSuccess);
 }
+
+// ================================================================================================
+hipError_t hipExtStreamCreateWithCUMask(hipStream_t* stream, uint32_t cuMaskSize,
+                                        const uint32_t* cuMask) {
+  HIP_INIT_API(hipExtStreamCreateWithCUMask, stream, cuMaskSize, cuMask);
+
+  if (stream == nullptr) {
+    HIP_RETURN(hipErrorInvalidHandle);
+  }
+  if (cuMaskSize == 0 || cuMask == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+
+  const std::vector<uint32_t> cuMaskv(cuMask, cuMask + cuMaskSize);
+
+  HIP_RETURN(ihipStreamCreate(stream, hipStreamDefault, amd::CommandQueue::Priority::Normal, cuMaskv));
+}
+

hipamd/tests/src/cg/hipCGThreadBlockType.cpp

@@ -0,0 +1,164 @@
+/*
+Copyright (c) 2020 - 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.
+*/
+
+
+/* HIT_START
+ * BUILD: %t %s ../test_common.cpp
+ * TEST: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+#include "hip/hip_cooperative_groups.h"
+
+#define HIP_ASSERT(lhs, rhs) assert(lhs == rhs)
+
+using namespace cooperative_groups;
+
+static __global__
+void kernel_cg_thread_block_type(dim3 *groupIndexD,
+                                 dim3 *thdIndexD,
+                                 int *sizeD,
+                                 int *thdRankD,
+                                 int *isValidD,
+                                 int *syncValD)
+{
+  thread_block tb = this_thread_block();
+
+  // Consider the workgroup id (0, 1, 1) and thread id (0, 1, 2) for validation
+  // of the test
+  int isBlockIdx011 =
+    (hipBlockIdx_x == 0 && hipBlockIdx_y == 1 && hipBlockIdx_z == 1);
+  int isThreadIdx012 =
+    (hipThreadIdx_x == 0 && hipThreadIdx_y == 1 && hipThreadIdx_z == 2);
+
+  if (isBlockIdx011 && isThreadIdx012) {
+    *groupIndexD = tb.group_index();
+    *thdIndexD = tb.thread_index();
+    *sizeD = tb.size();
+    *thdRankD = tb.thread_rank();
+    *isValidD = tb.is_valid();
+  }
+
+  // Consider local thread id (0, 0, 0) and (0, 0, 1) for validation of sync()
+  // api
+  __shared__ int sVar[2];
+  int isThreadIdx000 =
+    (hipThreadIdx_x == 0 && hipThreadIdx_y == 0 && hipThreadIdx_z == 0);
+  int isThreadIdx001 =
+    (hipThreadIdx_x == 0 && hipThreadIdx_y == 0 && hipThreadIdx_z == 1);
+
+  if (isThreadIdx000)
+    sVar[0] = 10;
+  if (isThreadIdx001)
+    sVar[1] = 20;
+  tb.sync();
+  if (isBlockIdx011 && isThreadIdx012)
+    *syncValD = sVar[0] + sVar[1];
+}
+
+static void test_cg_thread_block_type()
+{
+  dim3 *groupIndexD, *groupIndexH;
+  dim3 *thdIndexD, *thdIndexH;
+  int *sizeD, *sizeH;
+  int *thdRankD, *thdRankH;
+  int *isValidD, *isValidH;
+  int *syncValD, *syncValH;
+
+  // Allocate device memory
+  HIP_ASSERT(hipMalloc((void**)&groupIndexD, sizeof(dim3)), hipSuccess);
+  HIP_ASSERT(hipMalloc((void**)&thdIndexD, sizeof(dim3)), hipSuccess);
+  HIP_ASSERT(hipMalloc((void**)&sizeD, sizeof(int)), hipSuccess);
+  HIP_ASSERT(hipMalloc((void**)&thdRankD, sizeof(int)), hipSuccess);
+  HIP_ASSERT(hipMalloc((void**)&isValidD, sizeof(int)), hipSuccess);
+  HIP_ASSERT(hipMalloc((void**)&syncValD, sizeof(int)), hipSuccess);
+
+  // Allocate host memory
+  HIP_ASSERT(hipHostMalloc((void**)&groupIndexH, sizeof(dim3)), hipSuccess);
+  HIP_ASSERT(hipHostMalloc((void**)&thdIndexH, sizeof(dim3)), hipSuccess);
+  HIP_ASSERT(hipHostMalloc((void**)&sizeH, sizeof(int)), hipSuccess);
+  HIP_ASSERT(hipHostMalloc((void**)&thdRankH, sizeof(int)), hipSuccess);
+  HIP_ASSERT(hipHostMalloc((void**)&isValidH, sizeof(int)), hipSuccess);
+  HIP_ASSERT(hipHostMalloc((void**)&syncValH, sizeof(int)), hipSuccess);
+
+  // Launch Kernel
+  hipLaunchKernelGGL(kernel_cg_thread_block_type,
+                     dim3(2, 2, 2),
+                     dim3(4, 4, 4),
+                     0,
+                     0,
+                     groupIndexD,
+                     thdIndexD,
+                     sizeD,
+                     thdRankD,
+                     isValidD,
+                     syncValD);
+
+  // Copy result from device to host
+  HIP_ASSERT(hipMemcpy(groupIndexH, groupIndexD, sizeof(dim3), hipMemcpyDeviceToHost), hipSuccess);
+  HIP_ASSERT(hipMemcpy(thdIndexH, thdIndexD, sizeof(dim3), hipMemcpyDeviceToHost), hipSuccess);
+  HIP_ASSERT(hipMemcpy(sizeH, sizeD, sizeof(int), hipMemcpyDeviceToHost), hipSuccess);
+  HIP_ASSERT(hipMemcpy(thdRankH, thdRankD, sizeof(int), hipMemcpyDeviceToHost), hipSuccess);
+  HIP_ASSERT(hipMemcpy(isValidH, isValidD, sizeof(int), hipMemcpyDeviceToHost), hipSuccess);
+  HIP_ASSERT(hipMemcpy(syncValH, syncValD, sizeof(int), hipMemcpyDeviceToHost), hipSuccess);
+
+  // Validate result
+  // Group index should be (0, 1, 1)
+  HIP_ASSERT(groupIndexH->x, 0);
+  HIP_ASSERT(groupIndexH->y, 1);
+  HIP_ASSERT(groupIndexH->z, 1);
+  // Thread index should be (0, 1, 2)
+  HIP_ASSERT(thdIndexH->x, 0);
+  HIP_ASSERT(thdIndexH->y, 1);
+  HIP_ASSERT(thdIndexH->z, 2);
+  // Workgroup size should be 64
+  HIP_ASSERT(*sizeH, 64);
+  // Thread rank should be 36
+  HIP_ASSERT(*thdRankH, 36);
+  // Call to is_valid() should return true
+  HIP_ASSERT(*isValidH, 1);
+  // syncVal should be equal to 30
+  HIP_ASSERT(*syncValH, 30);
+
+  // Free device memory
+  HIP_ASSERT(hipFree(groupIndexD), hipSuccess);
+  HIP_ASSERT(hipFree(thdIndexD), hipSuccess);
+  HIP_ASSERT(hipFree(sizeD), hipSuccess);
+  HIP_ASSERT(hipFree(thdRankD), hipSuccess);
+  HIP_ASSERT(hipFree(isValidD), hipSuccess);
+  HIP_ASSERT(hipFree(syncValD), hipSuccess);
+
+  //Free host memory
+  HIP_ASSERT(hipHostFree(groupIndexH), hipSuccess);
+  HIP_ASSERT(hipHostFree(thdIndexH), hipSuccess);
+  HIP_ASSERT(hipHostFree(sizeH), hipSuccess);
+  HIP_ASSERT(hipHostFree(thdRankH), hipSuccess);
+  HIP_ASSERT(hipHostFree(isValidH), hipSuccess);
+  HIP_ASSERT(hipHostFree(syncValH), hipSuccess);
+}
+
+int main()
+{
+  test_cg_thread_block_type();
+  passed();
+}

hipamd/tests/src/texture/hipBindTex2DPitch.cpp

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

hipamd/tests/src/texture/hipBindTexRef1DFetch.cpp

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

hipamd/tests/src/texture/hipNormalizedFloatValueTex.cpp

@@ -21,7 +21,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc hcc rocclr
+ * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc hcc
  * TEST: %t
  * HIT_END
  */
@@ -42,24 +42,12 @@ static float getNormalizedValue(const float value,
     return value;
 }
 
-#if __HIP__
-__hip_pinned_shadow__
-#endif
 texture<char, hipTextureType1D, hipReadModeNormalizedFloat> texc;
 
-#if __HIP__
-__hip_pinned_shadow__
-#endif
 texture<unsigned char, hipTextureType1D, hipReadModeNormalizedFloat> texuc;
 
-#if __HIP__
-__hip_pinned_shadow__
-#endif
 texture<short, hipTextureType1D, hipReadModeNormalizedFloat> texs;
 
-#if __HIP__
-__hip_pinned_shadow__
-#endif
 texture<unsigned short, hipTextureType1D, hipReadModeNormalizedFloat> texus;
 
 

hipamd/tests/src/texture/hipTextureRef2D.cpp

@@ -1,5 +1,5 @@
 /* HIT_START
- * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM rocclr
+ * BUILD: %t %s ../test_common.cpp
  * TEST: %t
  * HIT_END
  */

hipamd/tests/src/texture/simpleTexture2DLayered.cpp

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

hipamd/tests/src/texture/simpleTexture3D.cpp

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