Merge 'master' into 'amd-master'
Change-Id: I44d884a9927410d3fafe0d43becdbd3819b544a2
Этот коммит содержится в:
Jenkins
@@ -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()
|
||||
|
||||
+75
-7
@@ -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__";
|
||||
}
|
||||
|
||||
@@ -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
|
||||
*/
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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.
|
||||
|
||||
@@ -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;
|
||||
|
||||
|
||||
@@ -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;
|
||||
|
||||
@@ -23,7 +23,7 @@ inline clara::Parser cmdline_parser(bool& help, std::vector<std::string>& 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<std::string>& inputs) {
|
||||
throw std::runtime_error{"Non existent file " + *it + " passed as input."};
|
||||
}
|
||||
}
|
||||
} // namespace hip_impl
|
||||
} // namespace hip_impl
|
||||
|
||||
@@ -12,7 +12,7 @@ namespace hip_impl {
|
||||
inline const std::unordered_set<std::string>& amdgpu_targets() { // The evolving list lives at:
|
||||
// https://www.llvm.org/docs/AMDGPUUsage.html#processors.
|
||||
static const std::unordered_set<std::string> r{"gfx701", "gfx801", "gfx802", "gfx803",
|
||||
"gfx900"};
|
||||
"gfx900", "gfx906"};
|
||||
|
||||
return r;
|
||||
}
|
||||
@@ -77,4 +77,4 @@ inline void validate_targets(const std::vector<std::string>& x) {
|
||||
}
|
||||
}
|
||||
}
|
||||
} // Namespace hip_impl.
|
||||
} // Namespace hip_impl.
|
||||
|
||||
@@ -132,9 +132,9 @@ inline clara::Parser cmdline_parser(bool& help, std::vector<std::string>& 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
|
||||
} // namespace hip_impl
|
||||
|
||||
@@ -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);
|
||||
}
|
||||
|
||||
/*
|
||||
|
||||
@@ -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<const __CudaFatBinaryWrapper*>(data);
|
||||
if (fbwrapper->magic != __cudaFatMAGIC2 || fbwrapper->version != 1) {
|
||||
if (fbwrapper->magic != __hipFatMAGIC2 || fbwrapper->version != 1) {
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
|
||||
+108
-21
@@ -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 <uint32_t block_dim, typename T>
|
||||
__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<void*>(dst+idx), reinterpret_cast<const void*>(src+idx),
|
||||
// sizeof(T));
|
||||
dst[idx] = src[idx];
|
||||
idx += grid_dim;
|
||||
}
|
||||
}
|
||||
} // namespace
|
||||
|
||||
template <typename T>
|
||||
@@ -1566,6 +1558,16 @@ void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t sizeBytes) {
|
||||
sizeBytes, std::move(val));
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
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<size_t>(sizeBytes / block_dim_, 1, UINT32_MAX);
|
||||
|
||||
hipLaunchKernelGGL(hip_copy_n<block_dim_>, 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<void*>(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<uint32_t> (stream, static_cast<uint32_t*> (dst), static_cast<const uint32_t*> (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);
|
||||
|
||||
|
||||
@@ -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;
|
||||
|
||||
@@ -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;
|
||||
|
||||
@@ -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 <assert.h>
|
||||
#include <stdio.h>
|
||||
#include <algorithm>
|
||||
#include <stdlib.h>
|
||||
#include <iostream>
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <hip/device_functions.h>
|
||||
|
||||
#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;
|
||||
}
|
||||
@@ -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<elements; i++) {
|
||||
A_h[i] = 1;
|
||||
}
|
||||
hipStream_t stream;
|
||||
//HIPCHECK(hipStreamCreate(&stream));
|
||||
HIPCHECK ( hipMemset3D( devPitchedPtr, memsetval, extent) );
|
||||
//HIPCHECK ( hipMemcpy3D(A_h, width, A_d, pitch_A, numW, numH, hipMemcpyDeviceToHost));
|
||||
hipMemcpy3DParms myparms = {0};
|
||||
myparms.srcPos = make_hipPos(0,0,0);
|
||||
myparms.dstPos = make_hipPos(0,0,0);
|
||||
myparms.dstPtr = make_hipPitchedPtr(A_h, width , numW, numH);
|
||||
myparms.srcPtr = devPitchedPtr;
|
||||
//myparms.dstArray = cuArray;
|
||||
myparms.extent = extent;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
|
||||
#else
|
||||
myparms.kind = hipMemcpyDeviceToHost;
|
||||
#endif
|
||||
hipMemcpy3D(&myparms);
|
||||
|
||||
for (int i=0; i<elements; 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;
|
||||
}
|
||||
}
|
||||
hipFree(A_d);
|
||||
free(A_h);
|
||||
return testResult;
|
||||
}
|
||||
|
||||
int main(int argc, char *argv[])
|
||||
{
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
bool testResult = false;
|
||||
HIPCHECK(hipSetDevice(p_gpuDevice));
|
||||
testResult = testhipMemset3D(memsetval, p_gpuDevice);
|
||||
|
||||
passed();
|
||||
|
||||
}
|
||||
Ссылка в новой задаче
Block a user