diff --git a/CMakeLists.txt b/CMakeLists.txt index 16c3f11edf..8278daad2c 100755 --- a/CMakeLists.txt +++ b/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") diff --git a/include/hip/hcc_detail/hip_cooperative_groups.h b/include/hip/hcc_detail/hip_cooperative_groups.h index 27ce887723..73f6778a1c 100644 --- a/include/hip/hcc_detail/hip_cooperative_groups.h +++ b/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(this)->thread_rank()); } + case internal::cg_workgroup: { + return (static_cast(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(this)->is_valid()); } + case internal::cg_workgroup: { + return (static_cast(this)->is_valid()); + } default: { return false; } @@ -207,6 +262,10 @@ __CG_QUALIFIER__ void thread_group::sync() const { static_cast(this)->sync(); break; } + case internal::cg_workgroup: { + static_cast(this)->sync(); + break; + } } } diff --git a/include/hip/hcc_detail/hip_cooperative_groups_helper.h b/include/hip/hcc_detail/hip_cooperative_groups_helper.h index 9738448d94..4e10c0da72 100644 --- a/include/hip/hcc_detail/hip_cooperative_groups_helper.h +++ b/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 diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 744885923f..d908f72a04 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/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 */ diff --git a/rocclr/CMakeLists.txt b/rocclr/CMakeLists.txt index 9d05de586d..37135af81d 100644 --- a/rocclr/CMakeLists.txt +++ b/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} $ $) @@ -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 - $ - ) - 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 $ # Must remove old one, otherwise the new one will contain obsolete stuff - COMMAND ${CMAKE_AR} -rcsT $ $ $ - 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 $ # Extract *.o from amdrocclr_static + COMMAND ${CMAKE_AR} -rcs $ *.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 $ DESTINATION lib COMPONENT MAIN) INSTALL(PROGRAMS $ DESTINATION lib COMPONENT MAIN) INSTALL(CODE "execute_process( COMMAND ${CMAKE_COMMAND} -E create_symlink libamdhip64.so lib/libhip_hcc.so )" DESTINATION lib COMPONENT MAIN) diff --git a/rocclr/hip_hcc.def.in b/rocclr/hip_hcc.def.in index 579608e685..53bdfb3065 100755 --- a/rocclr/hip_hcc.def.in +++ b/rocclr/hip_hcc.def.in @@ -252,3 +252,4 @@ hipTexObjectDestroy hipTexObjectGetResourceDesc hipTexObjectGetResourceViewDesc hipTexObjectGetTextureDesc +hipExtStreamCreateWithCUMask diff --git a/rocclr/hip_hcc.map.in b/rocclr/hip_hcc.map.in index 19da8a6991..02a704bf88 100755 --- a/rocclr/hip_hcc.map.in +++ b/rocclr/hip_hcc.map.in @@ -258,6 +258,7 @@ global: hipInitActivityCallback*; hipEnableActivityCallback*; hipGetCmdName*; + hipExtStreamCreateWithCUMask; }; local: *; diff --git a/rocclr/hip_internal.hpp b/rocclr/hip_internal.hpp index 4cc0dadd8a..492842c32b 100755 --- a/rocclr/hip_internal.hpp +++ b/rocclr/hip_internal.hpp @@ -87,9 +87,11 @@ namespace hip { amd::CommandQueue::Priority priority_; unsigned int flags_; bool null_; + const std::vector 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& cuMask = {}); /// Creates the hip stream object, including AMD host queue bool Create(); diff --git a/rocclr/hip_memory.cpp b/rocclr/hip_memory.cpp index 3aad7f57ba..b7b8d77592 100755 --- a/rocclr/hip_memory.cpp +++ b/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(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); } diff --git a/rocclr/hip_prof_gen.py b/rocclr/hip_prof_gen.py index c508888140..c59d1cca35 100755 --- a/rocclr/hip_prof_gen.py +++ b/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') diff --git a/rocclr/hip_stream.cpp b/rocclr/hip_stream.cpp index 3bd7d343f7..23b89e71ca 100644 --- a/rocclr/hip_stream.cpp +++ b/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& 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& 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 cuMaskv(cuMask, cuMask + cuMaskSize); + + HIP_RETURN(ihipStreamCreate(stream, hipStreamDefault, amd::CommandQueue::Priority::Normal, cuMaskv)); +} + diff --git a/tests/src/cg/hipCGThreadBlockType.cpp b/tests/src/cg/hipCGThreadBlockType.cpp new file mode 100644 index 0000000000..ab9492c609 --- /dev/null +++ b/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(); +} diff --git a/tests/src/texture/hipBindTex2DPitch.cpp b/tests/src/texture/hipBindTex2DPitch.cpp index 2fd3f1228d..3be39fb911 100644 --- a/tests/src/texture/hipBindTex2DPitch.cpp +++ b/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 */ diff --git a/tests/src/texture/hipBindTexRef1DFetch.cpp b/tests/src/texture/hipBindTexRef1DFetch.cpp index eeaf42129f..546e9775e8 100644 --- a/tests/src/texture/hipBindTexRef1DFetch.cpp +++ b/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 */ diff --git a/tests/src/texture/hipNormalizedFloatValueTex.cpp b/tests/src/texture/hipNormalizedFloatValueTex.cpp index 95b7c1d879..f365240543 100644 --- a/tests/src/texture/hipNormalizedFloatValueTex.cpp +++ b/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 texc; -#if __HIP__ -__hip_pinned_shadow__ -#endif texture texuc; -#if __HIP__ -__hip_pinned_shadow__ -#endif texture texs; -#if __HIP__ -__hip_pinned_shadow__ -#endif texture texus; diff --git a/tests/src/texture/hipTextureRef2D.cpp b/tests/src/texture/hipTextureRef2D.cpp index 4f00260998..e39c36638c 100644 --- a/tests/src/texture/hipTextureRef2D.cpp +++ b/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 */ diff --git a/tests/src/texture/simpleTexture2DLayered.cpp b/tests/src/texture/simpleTexture2DLayered.cpp index 1424976716..d8db1308e3 100644 --- a/tests/src/texture/simpleTexture2DLayered.cpp +++ b/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 */ diff --git a/tests/src/texture/simpleTexture3D.cpp b/tests/src/texture/simpleTexture3D.cpp index bd342df1c3..ca6695f619 100644 --- a/tests/src/texture/simpleTexture3D.cpp +++ b/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 */