diff --git a/CMakeLists.txt b/CMakeLists.txt index 075b916a4d..4f770b3429 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -201,7 +201,7 @@ if(HIP_PLATFORM STREQUAL "hcc") execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --ldflags OUTPUT_VARIABLE HCC_LD_FLAGS) set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} ${HCC_LD_FLAGS} -Wl,-Bsymbolic") - set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} --amdgpu-target=gfx701 --amdgpu-target=gfx803 --amdgpu-target=gfx900") + set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} --amdgpu-target=gfx701 --amdgpu-target=gfx803 --amdgpu-target=gfx900 --amdgpu-target=gfx906") if(COMPILE_HIP_ATP_MARKER) set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L/opt/rocm/profiler/CXLActivityLogger/bin/x86_64 -lCXLActivityLogger") endif() diff --git a/bin/hipcc b/bin/hipcc index f796b0bf95..b162de2dce 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -50,6 +50,8 @@ $verbose = $ENV{'HIPCC_VERBOSE'} // 0; # Verbose: 0x1=commands, 0x2=paths, 0x4=hipcc args $HIP_PATH=$ENV{'HIP_PATH'} // dirname (dirname $0); # use parent directory of hipcc +$HIP_CLANG_PATH=$ENV{'HIP_CLANG_PATH'}; +$DEVICE_LIB_PATH=$ENV{'DEVICE_LIB_PATH'}; #--- # Read .hipInfo @@ -62,6 +64,10 @@ $HIP_PLATFORM= `$HIP_PATH/bin/hipconfig --platform` // "hcc"; $HIP_VERSION= `$HIP_PATH/bin/hipconfig --version`; ($HIP_VERSION_MAJOR, $HIP_VERSION_MINOR, $HIP_VERSION_PATCH) = split(/\./, $HIP_VERSION); +if (defined $HIP_CLANG_PATH) { + $HIP_PLATFORM = "clang" +} + if ($verbose & 0x2) { print ("HIP_PATH=$HIP_PATH\n"); print ("HIP_PLATFORM=$HIP_PLATFORM\n"); @@ -75,9 +81,19 @@ $target_gfx801 = 0; $target_gfx802 = 0; $target_gfx803 = 0; $target_gfx900 = 0; +$target_gfx906 = 0; $default_amdgpu_target = 1; -if ($HIP_PLATFORM eq "hcc") { +if ($HIP_PLATFORM eq "clang") { + if ($verbose & 0x2) { + print ("HIP_CLANG_PATH=$HIP_CLANG_PATH\n"); + } + $ROCM_PATH=$ENV{'ROCM_PATH'} // "/opt/rocm"; + $HIPCC="$HIP_CLANG_PATH/clang++"; + $HIPCXXFLAGS .= " -I$HIP_PATH/include"; + + $HIPLDFLAGS = "--hip-link --hip-device-lib-path=$DEVICE_LIB_PATH -L$HIP_PATH/lib -lhip_hcc"; +} elsif ($HIP_PLATFORM eq "hcc") { $HSA_PATH=$ENV{'HSA_PATH'} // "/opt/rocm/hsa"; $HCC_HOME=$ENV{'HCC_HOME'} // $hipConfig{'HCC_HOME'} // "/opt/rocm/hcc"; @@ -281,6 +297,12 @@ foreach $arg (@ARGV) $target_gfx900 = 1; $default_amdgpu_target = 0; } + if($arg eq '--amdgpu-target=gfx906') + { + $target_gfx906 = 1; + $default_amdgpu_target = 0; + } + if(($trimarg eq '-stdlib=libstdc++') and ($setStdLib eq 0)) { @@ -330,10 +352,16 @@ foreach $arg (@ARGV) if (($arg =~ /\.cpp$/) or ($arg =~ /\.c$/) or ($arg =~ /\.cc$/) ) { $hasC = 1; $needCXXFLAGS = 1; + if ($HIP_PLATFORM eq 'clang') { + $toolArgs .= " -x hip" + } } if (($arg =~ /\.cu$/) or ($arg =~ /\.cuh$/)) { $hasCU = 1; $needCXXFLAGS = 1; + if ($HIP_PLATFORM eq 'clang') { + $toolArgs .= " -x hip" + } } push (@inputs, $arg); @@ -342,7 +370,7 @@ foreach $arg (@ARGV) $toolArgs .= " $arg" unless $swallowArg; } -if($HIP_PLATFORM eq "hcc"){ +if($HIP_PLATFORM eq "hcc" or $HIP_PLATFORM eq "clang"){ # No AMDGPU target specified at commandline. So look for HCC_AMDGPU_TARGET if($default_amdgpu_target eq 1 and defined $ENV{HCC_AMDGPU_TARGET}) { @@ -373,6 +401,11 @@ if($HIP_PLATFORM eq "hcc"){ $target_gfx900 = 1; $default_amdgpu_target = 0; } + if($target eq 'gfx906') + { + $target_gfx906 = 1; + $default_amdgpu_target = 0; + } } } # Else try using rocm_agent_enumerator @@ -404,6 +437,10 @@ if($HIP_PLATFORM eq "hcc"){ $target_gfx900 = 1; $default_amdgpu_target = 0; } + if($val eq "gfx906") { + $target_gfx906 = 1; + $default_amdgpu_target = 0; + } } } # rocm_agent_enumerator failed! Throw an error and die if linking is required @@ -414,29 +451,59 @@ if($HIP_PLATFORM eq "hcc"){ $ENV{HCC_EXTRA_LIBRARIES}="$HIP_PATH/lib/hip_hc.ll\n"; + if($HIP_PLATFORM eq "hcc") { + $GPU_ARCH_OPT = " --amdgpu-target="; + } else { + $GPU_ARCH_OPT = " --cuda-gpu-arch="; + } # Handle ROCm target platform if ($target_gfx701 eq 1) { - $HIPLDFLAGS .= " --amdgpu-target=gfx701"; + $GPU_ARCH_ARG = $GPU_ARCH_OPT . "gfx701"; + $HIPLDFLAGS .= $GPU_ARCH_ARG; + if ($HIP_PLATFORM eq 'clang') { + $HIPCXXFLAGS .= $GPU_ARCH_ARG;; + } $HIPCXXFLAGS .= " -D__HIP_ARCH_GFX701__=1 "; } if ($target_gfx801 eq 1) { - $HIPLDFLAGS .= " --amdgpu-target=gfx801"; + $GPU_ARCH_ARG = $GPU_ARCH_OPT . "gfx801"; + $HIPLDFLAGS .= $GPU_ARCH_ARG; + if ($HIP_PLATFORM eq 'clang') { + $HIPCXXFLAGS .= $GPU_ARCH_ARG;; + } $HIPCXXFLAGS .= " -D__HIP_ARCH_GFX801__=1 "; } if ($target_gfx802 eq 1) { - $HIPLDFLAGS .= " --amdgpu-target=gfx802"; + $GPU_ARCH_ARG = $GPU_ARCH_OPT . "gfx802"; + $HIPLDFLAGS .= $GPU_ARCH_ARG; + if ($HIP_PLATFORM eq 'clang') { + $HIPCXXFLAGS .= $GPU_ARCH_ARG;; + } $HIPCXXFLAGS .= " -D__HIP_ARCH_GFX802__=1 "; } if ($target_gfx803 eq 1) { - $HIPLDFLAGS .= " --amdgpu-target=gfx803"; + $GPU_ARCH_ARG = $GPU_ARCH_OPT . "gfx803"; + $HIPLDFLAGS .= $GPU_ARCH_ARG; + if ($HIP_PLATFORM eq 'clang') { + $HIPCXXFLAGS .= $GPU_ARCH_ARG;; + } $HIPCXXFLAGS .= " -D__HIP_ARCH_GFX803__=1 "; $ENV{HCC_EXTRA_LIBRARIES_GFX803}="$HIP_PATH/lib/hip_hc_gfx803.ll\n"; } if ($target_gfx900 eq 1) { - $HIPLDFLAGS .= " --amdgpu-target=gfx900"; + $GPU_ARCH_ARG = $GPU_ARCH_OPT . "gfx900"; + $HIPLDFLAGS .= $GPU_ARCH_ARG; + if ($HIP_PLATFORM eq 'clang') { + $HIPCXXFLAGS .= $GPU_ARCH_ARG;; + } $HIPCXXFLAGS .= " -D__HIP_ARCH_GFX900__=1 "; $ENV{HCC_EXTRA_LIBRARIES_GFX900}="$HIP_PATH/lib/hip_hc_gfx803.ll\n"; } + if ($target_gfx906 eq 1) { + $HIPLDFLAGS .= " --amdgpu-target=gfx906"; + $HIPCXXFLAGS .= " -D__HIP_ARCH_GFX906__=1 "; + $ENV{HCC_EXTRA_LIBRARIES_GFX906}="$HIP_PATH/lib/hip_hc_gfx803.ll\n"; + } } if ($hasC and $HIP_PLATFORM eq 'nvcc') { @@ -445,6 +512,7 @@ if ($hasC and $HIP_PLATFORM eq 'nvcc') { if ($hasCU and $HIP_PLATFORM eq 'hcc') { $HIPCXXFLAGS .= " -x c++"; } + if ($buildDeps and $HIP_PLATFORM eq 'nvcc') { $HIPCXXFLAGS .= " -M -D__CUDACC__"; } diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index 28d874b27a..804f34f4f6 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -345,6 +345,11 @@ __device__ static inline unsigned int __usad(unsigned int x, unsigned int y, uns return __hip_hc_ir_usad_int(x, y, z); } +extern __device__ __attribute__((const)) unsigned int __mbcnt_lo(unsigned int x, unsigned int y) __asm("llvm.amdgcn.mbcnt.lo"); +extern __device__ __attribute__((const)) unsigned int __mbcnt_hi(unsigned int x, unsigned int y) __asm("llvm.amdgcn.mbcnt.hi"); + +__device__ static inline unsigned int __lane_id() { return __mbcnt_hi(-1, __mbcnt_lo(-1, 0)); } + /* Rounding modes are not yet supported in HIP */ diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 92f06e9174..e1b334aec3 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -612,6 +612,22 @@ extern const __device__ __attribute__((weak)) __hip_builtin_gridDim_t gridDim; #define hipGridDim_y gridDim.y #define hipGridDim_z gridDim.z +#pragma push_macro("__DEVICE__") +#define __DEVICE__ extern "C" __device__ __attribute__((always_inline)) \ + __attribute__((weak)) + +__DEVICE__ void __device_trap() __asm("llvm.trap"); + +__DEVICE__ void inline __assert_fail(const char * __assertion, + const char *__file, + unsigned int __line, + const char *__function) +{ + // Ignore all the args for now. + __device_trap(); +} +#pragma push_macro("__DEVICE__") + #endif #endif // HIP_HCC_DETAIL_RUNTIME_H diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 8becab3c9c..ada0fac19a 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -1450,6 +1450,27 @@ hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, size_t height,hipStream_t stream __dparm(0)); +/** + * @brief Fills synchronously the memory area pointed to by pitchedDevPtr with the constant value. + * + * @param[in] pitchedDevPtr + * @param[in] value - constant value to be set + * @param[in] extent + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree + */ +hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ); + +/** + * @brief Fills asynchronously the memory area pointed to by pitchedDevPtr with the constant value. + * + * @param[in] pitchedDevPtr + * @param[in] value - constant value to be set + * @param[in] extent + * @param[in] stream + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree + */ +hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ,hipStream_t stream __dparm(0)); + /** * @brief Query memory info. * Return snapshot of free memory, and total allocatable memory on the device. diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index bfde2e942d..2b36f3e140 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -243,6 +243,8 @@ typedef enum __HIP_NODISCARD hipError_t { 1062, ///< Produced when trying to unlock a non-page-locked memory. hipErrorMapBufferObjectFailed = 1071, ///< Produced when the IPC memory attach failed from ROCr. + hipErrorAssert = + 1081, ///< Produced when the kernel calls assert. hipErrorTbd ///< Marker that more error codes are needed. } hipError_t; diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h index 6f222648a2..92393ab2e0 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/include/hip/nvcc_detail/hip_runtime_api.h @@ -160,6 +160,7 @@ typedef cudaSurfaceObject_t hipSurfaceObject_t; #define hipDeviceMapHost cudaDeviceMapHost #define hipExtent cudaExtent +#define hipPitchedPtr cudaPitchedPtr #define make_hipExtent make_cudaExtent #define make_hipPos make_cudaPos #define make_hipPitchedPtr make_cudaPitchedPtr @@ -649,6 +650,14 @@ inline static hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, si return hipCUDAErrorTohipError(cudaMemset2DAsync(dst, pitch, value, width, height, stream)); } +inline static hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ){ + return hipCUDAErrorTohipError(cudaMemset3D(pitchedDevPtr, value, extent)); +} + +inline static hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent, hipStream_t stream __dparm(0) ){ + return hipCUDAErrorTohipError(cudaMemset3DAsync(pitchedDevPtr, value, extent, stream)); +} + inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t* p_prop, int device) { struct cudaDeviceProp cdprop; cudaError_t cerror; diff --git a/lpl_ca/ca.hpp b/lpl_ca/ca.hpp index bb1963bede..0ef8458c20 100644 --- a/lpl_ca/ca.hpp +++ b/lpl_ca/ca.hpp @@ -23,7 +23,7 @@ inline clara::Parser cmdline_parser(bool& help, std::vector& inputs "https://reviews.llvm.org/D13909; " "the code object format is documented at: " "https://www.llvm.org/docs/AMDGPUUsage.html#code-object.") | - clara::Opt{targets, "gfx803,gfx900 etc."}["-t"]["--targets"]( + clara::Opt{targets, "gfx803,gfx900,gfx906 etc."}["-t"]["--targets"]( "targets for which code objects are to be extracted from " "the fat binary; must be included in the set of processors " "with ROCm support from " @@ -76,4 +76,4 @@ inline void validate_inputs(const std::vector& inputs) { throw std::runtime_error{"Non existent file " + *it + " passed as input."}; } } -} // namespace hip_impl \ No newline at end of file +} // namespace hip_impl diff --git a/lpl_ca/common.hpp b/lpl_ca/common.hpp index 74f4db9221..7652f08c46 100644 --- a/lpl_ca/common.hpp +++ b/lpl_ca/common.hpp @@ -12,7 +12,7 @@ namespace hip_impl { inline const std::unordered_set& amdgpu_targets() { // The evolving list lives at: // https://www.llvm.org/docs/AMDGPUUsage.html#processors. static const std::unordered_set r{"gfx701", "gfx801", "gfx802", "gfx803", - "gfx900"}; + "gfx900", "gfx906"}; return r; } @@ -77,4 +77,4 @@ inline void validate_targets(const std::vector& x) { } } } -} // Namespace hip_impl. \ No newline at end of file +} // Namespace hip_impl. diff --git a/lpl_ca/lpl.hpp b/lpl_ca/lpl.hpp index c9346fd793..cbd7fe8386 100644 --- a/lpl_ca/lpl.hpp +++ b/lpl_ca/lpl.hpp @@ -132,9 +132,9 @@ inline clara::Parser cmdline_parser(bool& help, std::vector& source "file is documented at: https://reviews.llvm.org/D13909.") | clara::Arg{sources, "a.cpp b.cpp etc."}("inputs for compilation; must contain valid C++ code.") | - clara::Opt{targets, "gfx803,gfx900 etc."}["-t"]["--targets"]( + clara::Opt{targets, "gfx803,gfx900,gfx906 etc."}["-t"]["--targets"]( "targets for AMDGPU lowering; must be included in the set " "of processors with ROCm support from " "https://www.llvm.org/docs/AMDGPUUsage.html#processors."); } -} // namespace hip_impl \ No newline at end of file +} // namespace hip_impl diff --git a/src/device_functions.cpp b/src/device_functions.cpp index 86d0530817..879d354337 100644 --- a/src/device_functions.cpp +++ b/src/device_functions.cpp @@ -283,21 +283,6 @@ struct uchar2Holder { }; } __attribute__((aligned(8))); -struct intHolder { - union { - signed int si[2]; - signed int long sl; - }; -} __attribute__((aligned(8))); - -struct uintHolder { - union { - signed int ui[2]; - signed int long ul; - }; -} __attribute__((aligned(8))); - - __device__ unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s) { struct uchar2Holder cHoldVal; struct ucharHolder cHoldKey; @@ -313,21 +298,29 @@ __device__ unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int } __device__ long long __mul64hi(long long int x, long long int y) { - struct intHolder iHold1; - struct intHolder iHold2; - iHold1.sl = x; - iHold2.sl = y; - iHold1.sl = iHold1.si[1] * iHold2.si[1]; - return iHold1.sl; + ulong x0 = (ulong)x & 0xffffffffUL; + long x1 = x >> 32; + ulong y0 = (ulong)y & 0xffffffffUL; + long y1 = y >> 32; + ulong z0 = x0*y0; + long t = x1*y0 + (z0 >> 32); + long z1 = t & 0xffffffffL; + long z2 = t >> 32; + z1 = x0*y1 + z1; + return x1*y1 + z2 + (z1 >> 32); } __device__ unsigned long long __umul64hi(unsigned long long int x, unsigned long long int y) { - struct uintHolder uHold1; - struct uintHolder uHold2; - uHold1.ul = x; - uHold2.ul = y; - uHold1.ul = uHold1.ui[1] * uHold2.ui[1]; - return uHold1.ul; + ulong x0 = x & 0xffffffffUL; + ulong x1 = x >> 32; + ulong y0 = y & 0xffffffffUL; + ulong y1 = y >> 32; + ulong z0 = x0*y0; + ulong t = x1*y0 + (z0 >> 32); + ulong z1 = t & 0xffffffffUL; + ulong z2 = t >> 32; + z1 = x0*y1 + z1; + return x1*y1 + z2 + (z1 >> 32); } /* diff --git a/src/hip_clang.cpp b/src/hip_clang.cpp index 80b6111fc2..f7573e2819 100644 --- a/src/hip_clang.cpp +++ b/src/hip_clang.cpp @@ -27,10 +27,10 @@ THE SOFTWARE. #include "hip_hcc_internal.h" #include "trace_helper.h" -constexpr unsigned __cudaFatMAGIC2 = 0x466243b1; +constexpr unsigned __hipFatMAGIC2 = 0x48495046; // "HIPF" #define CLANG_OFFLOAD_BUNDLER_MAGIC "__CLANG_OFFLOAD_BUNDLE__" -#define AMDGCN_AMDHSA_TRIPLE "openmp-amdgcn--amdhsa" +#define AMDGCN_AMDHSA_TRIPLE "hip-amdgcn-amd-amdhsa" struct __ClangOffloadBundleDesc { uint64_t offset; @@ -59,7 +59,7 @@ __hipRegisterFatBinary(const void* data) HIP_INIT(); const __CudaFatBinaryWrapper* fbwrapper = reinterpret_cast(data); - if (fbwrapper->magic != __cudaFatMAGIC2 || fbwrapper->version != 1) { + if (fbwrapper->magic != __hipFatMAGIC2 || fbwrapper->version != 1) { return nullptr; } diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index cd744c8cc0..e2202e7860 100644 --- a/src/hip_memory.cpp +++ b/src/hip_memory.cpp @@ -1303,27 +1303,6 @@ hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) { return ihipLogStatus(e); } -hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, - size_t height, hipMemcpyKind kind, hipStream_t stream) { - HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind, stream); - if (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return ihipLogStatus(hipErrorInvalidValue); - hipError_t e = hipSuccess; - if((width == dpitch) && (width == spitch)) { - hip_internal::memcpyAsync(dst, src, width*height, kind, stream); - } else { - try { - for (int i = 0; i < height; ++i) { - e = hip_internal::memcpyAsync((unsigned char*)dst + i * dpitch, - (unsigned char*)src + i * spitch, width, kind, stream); - } - } catch (ihipException& ex) { - e = ex._code; - } - } - - return ihipLogStatus(e); -} - hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind) { HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, wOffset, hOffset, src, spitch, width, height, kind); @@ -1554,6 +1533,19 @@ inline const T& clamp_integer(const T& x, const T& lower, const T& upper) { return std::min(upper, std::max(x, lower)); } + +template +__global__ void hip_copy_n(T* dst, const T* src, size_t n) { + const uint32_t grid_dim = gridDim.x * blockDim.x; + + size_t idx = blockIdx.x * block_dim + threadIdx.x; + while (idx < n) { + // __builtin_memcpy(reinterpret_cast(dst+idx), reinterpret_cast(src+idx), + // sizeof(T)); + dst[idx] = src[idx]; + idx += grid_dim; + } +} } // namespace template @@ -1566,6 +1558,16 @@ void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t sizeBytes) { sizeBytes, std::move(val)); } +template +void ihipMemcpyKernel(hipStream_t stream, T* dst, const T* src, size_t sizeBytes) { + static constexpr uint32_t block_dim_ = 256; + + const uint32_t grid_dim = clamp_integer(sizeBytes / block_dim_, 1, UINT32_MAX); + + hipLaunchKernelGGL(hip_copy_n, dim3(grid_dim), dim3{block_dim_}, 0u, stream, dst, src, + sizeBytes); +} + typedef enum ihipMemsetDataType { ihipMemsetDataTypeChar = 0, ihipMemsetDataTypeShort = 1, @@ -1623,6 +1625,55 @@ hipError_t ihipMemset(void* dst, int value, size_t sizeBytes, hipStream_t strea return e; }; +int isLockedPointer(const void *ptr) +{ + hsa_amd_pointer_info_t info; + int isLocked = 0; + + info.size = sizeof(info); + hsa_status_t hsa_status = hsa_amd_pointer_info(const_cast(ptr), &info, nullptr, nullptr, nullptr); + if(hsa_status != HSA_STATUS_SUCCESS) { + return -1; + } + + if((info.type == HSA_EXT_POINTER_TYPE_HSA) || (info.type == HSA_EXT_POINTER_TYPE_LOCKED)) { + isLocked = 1; + } + + return isLocked; +} + +hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, + size_t height, hipMemcpyKind kind, hipStream_t stream) { + HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind, stream); + if (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return ihipLogStatus(hipErrorInvalidValue); + hipError_t e = hipSuccess; + int isLocked = 0; + if(kind == hipMemcpyHostToDevice) { + isLocked = isLockedPointer(src); + } else if(kind == hipMemcpyDeviceToHost) { + isLocked = isLockedPointer(dst); + } + if((width == dpitch) && (width == spitch)) { + hip_internal::memcpyAsync(dst, src, width*height, kind, stream); + } else { + try { + for (int i = 0; i < height; ++i) { + if(!isLocked) { + e = hip_internal::memcpyAsync((unsigned char*)dst + i * dpitch, + (unsigned char*)src + i * spitch, width, kind, stream); + } else{ + size_t sizeBytes = width*height; + ihipMemcpyKernel (stream, static_cast (dst), static_cast (src), sizeBytes/sizeof(uint32_t)); + } + } + } catch (ihipException& ex) { + e = ex._code; + } + } + + return ihipLogStatus(e); +} // TODO-sync: function is async unless target is pinned host memory - then these are fully sync. hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) { @@ -1705,6 +1756,42 @@ hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes return ihipLogStatus(e); } +hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ) +{ + HIP_INIT_SPECIAL_API((TRACE_MCMD), &pitchedDevPtr, value, &extent); + hipError_t e = hipSuccess; + + hipStream_t stream = hipStreamNull; + // TODO - call an ihip memset so HIP_TRACE is correct. + stream = ihipSyncAndResolveStream(stream); + if (stream) { + size_t sizeBytes = pitchedDevPtr.pitch * extent.height * extent.depth; + e = ihipMemset(pitchedDevPtr.ptr, value, sizeBytes, stream, ihipMemsetDataTypeChar); + stream->locked_wait(); + } else { + e = hipErrorInvalidValue; + } + + return ihipLogStatus(e); +} + +hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ,hipStream_t stream ) +{ + HIP_INIT_SPECIAL_API((TRACE_MCMD), &pitchedDevPtr, value, &extent); + hipError_t e = hipSuccess; + + // TODO - call an ihip memset so HIP_TRACE is correct. + stream = ihipSyncAndResolveStream(stream); + if (stream) { + size_t sizeBytes = pitchedDevPtr.pitch * extent.height * extent.depth; + e = ihipMemset(pitchedDevPtr.ptr, value, sizeBytes, stream, ihipMemsetDataTypeChar); + } else { + e = hipErrorInvalidValue; + } + + return ihipLogStatus(e); +} + hipError_t hipMemGetInfo(size_t* free, size_t* total) { HIP_INIT_API(free, total); diff --git a/tests/src/deviceLib/hipMathFunctions.cpp b/tests/src/deviceLib/hipMathFunctions.cpp index 7fe0003672..78e85ba62a 100644 --- a/tests/src/deviceLib/hipMathFunctions.cpp +++ b/tests/src/deviceLib/hipMathFunctions.cpp @@ -29,7 +29,7 @@ THE SOFTWARE. #include "hip/hip_runtime.h" #include "test_common.h" -#if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX900__ +#if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX900__ || __HIP_ARCH_GFX906__ __global__ void kernel_abs_int64(hipLaunchParm lp, long long *input, long long *output) { int tx = threadIdx.x; diff --git a/tests/src/deviceLib/hipTestHalf.cpp b/tests/src/deviceLib/hipTestHalf.cpp index 4a63260bf7..24a4d6c53e 100644 --- a/tests/src/deviceLib/hipTestHalf.cpp +++ b/tests/src/deviceLib/hipTestHalf.cpp @@ -32,7 +32,7 @@ THE SOFTWARE. #define HALF_SIZE 64 * sizeof(__half) #define HALF2_SIZE 64 * sizeof(__half2) -#if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX900__ +#if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX900__ || __HIP_ARCH_GFX906__ __global__ void __halfMath(hipLaunchParm lp, __half* A, __half* B, __half* C) { int tx = threadIdx.x; diff --git a/tests/src/deviceLib/hip_mbcnt.cpp b/tests/src/deviceLib/hip_mbcnt.cpp new file mode 100644 index 0000000000..0dd7169f51 --- /dev/null +++ b/tests/src/deviceLib/hip_mbcnt.cpp @@ -0,0 +1,128 @@ +/* +Copyright (c) 2015-2018 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 EXCLUDE_HIP_PLATFORM nvcc + * RUN: %t + * HIT_END + */ + +#include +#include +#include +#include +#include +#include +#include + +#define HIP_ASSERT(x) (assert((x) == hipSuccess)) + +__global__ void HIP_kernel(hipLaunchParm lp, unsigned int* mbcnt_lo, unsigned int* mbcnt_hi, unsigned int* lane_id) { + int x = blockDim.x * blockIdx.x + threadIdx.x; + mbcnt_lo[x] = __mbcnt_lo(0xFFFFFFFF, 0); + mbcnt_hi[x] = __mbcnt_hi(0xFFFFFFFF, 0); + lane_id[x] = __lane_id(); +} + +using namespace std; + +int main() { + + unsigned int* device_mbcnt_lo; + unsigned int* device_mbcnt_hi; + unsigned int* device_lane_id; + + hipDeviceProp_t devProp; + hipGetDeviceProperties(&devProp, 0); + cout << " System minor " << devProp.minor << endl; + cout << " System major " << devProp.major << endl; + cout << " agent prop name " << devProp.name << endl; + + cout << "hip Device prop succeeded " << endl; + + constexpr unsigned int wave_size = 64; + constexpr unsigned int num_waves_per_block = 2; + constexpr unsigned int num_threads_per_block = wave_size * num_waves_per_block; + constexpr unsigned int num_blocks = 2; + constexpr unsigned int num_threads = num_threads_per_block * num_blocks; + constexpr size_t buffer_size = num_threads * sizeof(unsigned int); + + HIP_ASSERT(hipMalloc((void**)&device_mbcnt_lo, buffer_size)); + HIP_ASSERT(hipMalloc((void**)&device_mbcnt_hi, buffer_size)); + HIP_ASSERT(hipMalloc((void**)&device_lane_id, buffer_size)); + + hipLaunchKernel(HIP_kernel, dim3(num_blocks), + dim3(num_threads_per_block), 0, 0, device_mbcnt_lo, device_mbcnt_hi, device_lane_id); + + unsigned int* host_mbcnt_lo = (unsigned int*) malloc(buffer_size); + unsigned int* host_mbcnt_hi = (unsigned int*) malloc(buffer_size); + unsigned int* host_lane_id = (unsigned int*) malloc(buffer_size); + + HIP_ASSERT(hipMemcpy(host_mbcnt_lo, device_mbcnt_lo, buffer_size, hipMemcpyDeviceToHost)); + HIP_ASSERT(hipMemcpy(host_mbcnt_hi, device_mbcnt_hi, buffer_size, hipMemcpyDeviceToHost)); + HIP_ASSERT(hipMemcpy(host_lane_id, device_lane_id, buffer_size, hipMemcpyDeviceToHost)); + + // verify the results + int mbcnt_lo_errors = 0; + int mbcnt_hi_errors = 0; + int lane_id_errors = 0; + for (unsigned int i = 0; i < num_threads; i++) { + unsigned int this_lane_id = i % wave_size; + unsigned int this_mbcnt_lo = this_lane_id >= 32 ? 32 : this_lane_id; + unsigned int this_mbcnt_hi = this_lane_id < 32 ? 0 : (this_lane_id - 22); + + if (host_mbcnt_lo[i] != this_mbcnt_lo) + mbcnt_lo_errors++; + + if (host_mbcnt_hi[i] != this_mbcnt_hi) + mbcnt_hi_errors++; + + if (host_lane_id[i] != this_lane_id) + lane_id_errors++; + } + + if (mbcnt_lo_errors == 0) + cout << "__mbcnt_lo() PASSED!" << endl; + else + cout << "__mbcnt_lo() FAILED!" << endl; + + + if (mbcnt_hi_errors == 0) + cout << "__mbcnt_hi() PASSED!" << endl; + else + cout << "__mbcnt_hi() FAILED!" << endl; + + if (lane_id_errors == 0) + cout << "__lane_id() PASSED!" << endl; + else + cout << "__lane_id() FAILED!" << endl; + + HIP_ASSERT(hipFree(device_mbcnt_lo)); + HIP_ASSERT(hipFree(device_mbcnt_hi)); + HIP_ASSERT(hipFree(device_lane_id)); + + free(host_mbcnt_lo); + free(host_mbcnt_hi); + free(host_lane_id); + + return mbcnt_lo_errors + mbcnt_hi_errors + lane_id_errors; +} diff --git a/tests/src/runtimeApi/memory/hipMemset3D.cpp b/tests/src/runtimeApi/memory/hipMemset3D.cpp new file mode 100644 index 0000000000..40f2f3e67f --- /dev/null +++ b/tests/src/runtimeApi/memory/hipMemset3D.cpp @@ -0,0 +1,98 @@ +/* +Copyright (c) 2015-2016 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. +*/ +// Simple test for memset. +// Also serves as a template for other tests. + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp + * //Small copy + * RUN: %t -N 10 --memsetval 0x42 + * HIT_END + */ + +#include "hip/hip_runtime.h" +#include "test_common.h" + +bool testhipMemset3D(int memsetval,int p_gpuDevice) +{ + size_t numH = 256; + size_t numW = 256; + size_t depth = 1; + size_t pitch_A; + size_t width = numW * sizeof(char); + size_t sizeElements = width * numH * depth; + size_t elements = numW* numH* depth; + + + printf ("testhipMemset3D memsetval=%2x device=%d\n", memsetval, p_gpuDevice); + char *A_d; + char *A_h; + bool testResult = true; + hipExtent extent = make_hipExtent(width, numH, depth); + hipPitchedPtr devPitchedPtr; + + HIPCHECK(hipMalloc3D(&devPitchedPtr, extent)); + A_h = (char*)malloc(sizeElements); + HIPASSERT(A_h != NULL); + for (size_t i=0; i