AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge branch 'amd-master' into amd-master-next

Browse Source 
Change-Id: I3094c15008093f2072bcd38aca4ea90aeae2d97b
This commit is contained in:
Maneesh Gupta
2020-04-07 06:57:42 -04:00
parent 35b001b33a
commit 0ea6697192
67 changed files with 2278 additions and 601 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/include/hip/hcc_detail/code_object_bundle.hpp
-193
View File
@@ -1,193 +0,0 @@
/*
Copyright (c) 2015 - 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.
*/
#pragma once
#include <hsa/hsa.h>
#include <algorithm>
#include <cstdint>
#include <istream>
#include <iterator>
#include <string>
#include <utility>
#include <vector>
#include <unordered_set>
namespace hip_impl {
#if !defined(DISABLE_REDUCED_GPU_BLOB_COPY)
std::unordered_set<std::string>& get_all_gpuarch();
#endif
inline
std::string transmogrify_triple(const std::string& triple)
{
static constexpr const char old_prefix[]{"hcc-amdgcn--amdhsa-gfx"};
static constexpr const char new_prefix[]{"hcc-amdgcn-amd-amdhsa--gfx"};
if (triple.find(old_prefix) == 0) {
return new_prefix + triple.substr(sizeof(old_prefix) - 1);
}
return (triple.find(new_prefix) == 0) ? triple : "";
}
inline
std::string isa_name(std::string triple)
{
static constexpr const char offload_prefix[]{"hcc-"};
triple = transmogrify_triple(triple);
if (triple.empty()) return {};
triple.erase(0, sizeof(offload_prefix) - 1);
return triple;
}
inline
hsa_isa_t triple_to_hsa_isa(const std::string& triple) {
const std::string isa{isa_name(std::move(triple))};
if (isa.empty()) return hsa_isa_t({});
hsa_isa_t r{};
if(HSA_STATUS_SUCCESS != hsa_isa_from_name(isa.c_str(), &r)) {
r.handle = 0;
}
return r;
}
struct Bundled_code {
union Header {
struct {
std::uint64_t offset;
std::uint64_t bundle_sz;
std::uint64_t triple_sz;
};
char cbuf[sizeof(offset) + sizeof(bundle_sz) + sizeof(triple_sz)];
} header;
std::string triple;
std::string blob;
};
#define magic_string_ "__CLANG_OFFLOAD_BUNDLE__"
#ifdef __GNUC__
#pragma GCC visibility push (default)
#endif
class Bundled_code_header {
// DATA - STATICS
static constexpr auto magic_string_sz_ = sizeof(magic_string_) - 1;
// DATA
union Header_ {
struct {
char bundler_magic_string_[magic_string_sz_];
std::uint64_t bundle_cnt_;
};
char cbuf_[sizeof(bundler_magic_string_) + sizeof(bundle_cnt_)];
} header_;
std::vector<Bundled_code> bundles_;
// FRIENDS - MANIPULATORS
template <typename RandomAccessIterator>
friend inline bool read(RandomAccessIterator f, RandomAccessIterator l,
Bundled_code_header& x) {
if (f == l) return false;
std::copy_n(f, sizeof(x.header_.cbuf_), x.header_.cbuf_);
if (valid(x)) {
x.bundles_.resize(x.header_.bundle_cnt_);
auto it = f + sizeof(x.header_.cbuf_);
for (auto&& y : x.bundles_) {
std::copy_n(it, sizeof(y.header.cbuf), y.header.cbuf);
it += sizeof(y.header.cbuf);
y.triple.assign(it, it + y.header.triple_sz);
#ifdef DISABLE_REDUCED_GPU_BLOB_COPY
std::copy_n(f + y.header.offset, y.header.bundle_sz, std::back_inserter(y.blob));
#else
auto& gpuArch = get_all_gpuarch();
auto itgpuArch = std::find(gpuArch.begin(),gpuArch.end(),y.triple);
if (itgpuArch != gpuArch.end()){
std::copy_n(f + y.header.offset, y.header.bundle_sz, std::back_inserter(y.blob));
}
#endif
it += y.header.triple_sz;
x.bundled_code_size = std::max(x.bundled_code_size,
y.header.offset + y.header.bundle_sz);
}
return true;
}
return false;
}
friend inline bool read(const std::vector<char>& blob, Bundled_code_header& x) {
return read(blob.cbegin(), blob.cend(), x);
}
friend inline bool read(std::istream& is, Bundled_code_header& x) {
return read(
std::vector<char>{std::istreambuf_iterator<char>{is}, std::istreambuf_iterator<char>{}},
x);
}
// FRIENDS - ACCESSORS
friend inline bool valid(const Bundled_code_header& x) {
const std::string ms = {magic_string_};
return std::equal(ms.begin(), ms.end(), x.header_.bundler_magic_string_);
}
friend inline const std::vector<Bundled_code>& bundles(const Bundled_code_header& x) {
return x.bundles_;
}
public:
// CREATORS
Bundled_code_header() = default;
template <typename RandomAccessIterator>
Bundled_code_header(RandomAccessIterator f, RandomAccessIterator l);
explicit Bundled_code_header(const std::vector<char>& blob);
explicit Bundled_code_header(const void* maybe_blob);
Bundled_code_header(const Bundled_code_header&) = default;
Bundled_code_header(Bundled_code_header&&) = default;
~Bundled_code_header() = default;
// MANIPULATORS
Bundled_code_header& operator=(const Bundled_code_header&) = default;
Bundled_code_header& operator=(Bundled_code_header&&) = default;
size_t bundled_code_size = 0;
};
#ifdef __GNUC__
#pragma GCC visibility pop
#endif
// CREATORS
template <typename RandomAccessIterator>
Bundled_code_header::Bundled_code_header(RandomAccessIterator f, RandomAccessIterator l)
: Bundled_code_header{} {
read(f, l, *this);
}
} // Namespace hip_impl.
hipamd/include/hip/hcc_detail/device_functions.h
+4 -4
View File
@@ -128,7 +128,7 @@ __device__ static int __mul24(int x, int y);
__device__ static long long int __mul64hi(long long int x, long long int y);
__device__ static int __mulhi(int x, int y);
__device__ static int __rhadd(int x, int y);
__device__ static unsigned int __sad(int x, int y, int z);
__device__ static unsigned int __sad(int x, int y,unsigned int z);
__device__ static unsigned int __uhadd(unsigned int x, unsigned int y);
__device__ static int __umul24(unsigned int x, unsigned int y);
__device__ static unsigned long long int __umul64hi(unsigned long long int x, unsigned long long int y);
@@ -199,7 +199,7 @@ __device__ static inline int __rhadd(int x, int y) {
int value = z & 0x7FFFFFFF;
return ((value) >> 1 || sign);
}
__device__ static inline unsigned int __sad(int x, int y, int z) {
__device__ static inline unsigned int __sad(int x, int y, unsigned int z) {
return x > y ? x - y + z : y - x + z;
}
__device__ static inline unsigned int __uhadd(unsigned int x, unsigned int y) {
@@ -230,7 +230,7 @@ __device__ static inline unsigned int __urhadd(unsigned int x, unsigned int y) {
return (x + y + 1) >> 1;
}
__device__ static inline unsigned int __usad(unsigned int x, unsigned int y, unsigned int z) {
return __ockl_sad_u32(x, y, z);
return __ockl_sadd_u32(x, y, z);
}
__device__ static inline unsigned int __lane_id() { return __mbcnt_hi(-1, __mbcnt_lo(-1, 0)); }
@@ -563,7 +563,7 @@ long __shfl_xor(long var, int lane_mask, int width = warpSize)
return tmp1;
#else
static_assert(sizeof(long) == sizeof(int), "");
return static_cast<long>(__shfl_down(static_cast<int>(var), lane_mask, width));
return static_cast<long>(__shfl_xor(static_cast<int>(var), lane_mask, width));
#endif
}
__device__
hipamd/include/hip/hcc_detail/device_library_decls.h
+2 -1
View File
@@ -44,7 +44,7 @@ extern "C" __device__ __attribute__((const)) uint __ockl_mul24_u32(uint, uint);
extern "C" __device__ __attribute__((const)) int __ockl_mul24_i32(int, int);
extern "C" __device__ __attribute__((const)) uint __ockl_mul_hi_u32(uint, uint);
extern "C" __device__ __attribute__((const)) int __ockl_mul_hi_i32(int, int);
extern "C" __device__ __attribute__((const)) uint __ockl_sad_u32(uint, uint, uint);
extern "C" __device__ __attribute__((const)) uint __ockl_sadd_u32(uint, uint, uint);
extern "C" __device__ __attribute__((const)) uchar __ockl_clz_u8(uchar);
extern "C" __device__ __attribute__((const)) ushort __ockl_clz_u16(ushort);
@@ -72,6 +72,7 @@ extern "C" __device__ __attribute__((const)) uint __ockl_multi_grid_thread_rank(
extern "C" __device__ __attribute__((const)) int __ockl_multi_grid_is_valid(void);
extern "C" __device__ __attribute__((convergent)) void __ockl_multi_grid_sync(void);
extern "C" __device__ void __ockl_atomic_add_noret_f32(float*, float);
// Introduce local address space
#define __local __attribute__((address_space(3)))
hipamd/include/hip/hcc_detail/functional_grid_launch.hpp
+17 -14
View File
@@ -37,14 +37,15 @@ THE SOFTWARE.
hipError_t ihipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices,
unsigned int flags, hip_impl::program_state& ps);
hipError_t ihipLaunchCooperativeKernel(const void* f, dim3 gridDim, dim3 blockDimX, void** kernelParams,
unsigned int sharedMemBytes, hipStream_t stream, hip_impl::program_state& ps);
hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices,
unsigned int flags, hip_impl::program_state& ps);
hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim,
dim3 blockDim, void** args,
size_t sharedMem, hipStream_t stream,
hip_impl::program_state& ps);
hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList,
int numDevices,
unsigned int flags,
hip_impl::program_state& ps);
#pragma GCC visibility push(hidden)
@@ -192,22 +193,24 @@ void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
template <typename F>
inline
__attribute__((visibility("hidden")))
hipError_t hipLaunchCooperativeKernel(F f, dim3 gridDim, dim3 blockDimX, void** kernelParams,
unsigned int sharedMemBytes, hipStream_t stream) {
hipError_t hipLaunchCooperativeKernel(F f, dim3 gridDim, dim3 blockDim,
void** args, size_t sharedMem,
hipStream_t stream) {
hip_impl::hip_init();
auto& ps = hip_impl::get_program_state();
return ihipLaunchCooperativeKernel(reinterpret_cast<void*>(f), gridDim, blockDimX, kernelParams, sharedMemBytes, stream, ps);
return hipLaunchCooperativeKernel(reinterpret_cast<void*>(f), gridDim,
blockDim, args, sharedMem, stream, ps);
}
inline
__attribute__((visibility("hidden")))
hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices,
unsigned int flags) {
hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList,
int numDevices,
unsigned int flags) {
hip_impl::hip_init();
auto& ps = hip_impl::get_program_state();
return ihipLaunchCooperativeKernelMultiDevice(launchParamsList, numDevices, flags, ps);
return hipLaunchCooperativeKernelMultiDevice(launchParamsList, numDevices, flags, ps);
}
#pragma GCC visibility pop
hipamd/include/hip/hcc_detail/hip_atomic.h
+8
View File
@@ -73,6 +73,14 @@ float atomicAdd(float* address, float val)
return __uint_as_float(r);
}
__device__
inline
void atomicAddNoRet(float* address, float val)
{
__ockl_atomic_add_noret_f32(address, val);
}
__device__
inline
double atomicAdd(double* address, double val)
hipamd/include/hip/hcc_detail/hip_fp16.h
+14
View File
@@ -1268,6 +1268,13 @@ THE SOFTWARE.
static_cast<__half_raw>(x).data +
static_cast<__half_raw>(y).data};
}
inline
__device__
__half __habs(__half x)
{
return __half_raw{
__ocml_fabs_f16(static_cast<__half_raw>(x).data)};
}
inline
__device__
__half __hsub(__half x, __half y)
@@ -1334,6 +1341,13 @@ THE SOFTWARE.
static_cast<__half2_raw>(x).data +
static_cast<__half2_raw>(y).data};
}
inline
__device__
__half2 __habs2(__half2 x)
{
return __half2_raw{
__ocml_fabs_2f16(static_cast<__half2_raw>(x).data)};
}
inline
__device__
__half2 __hsub2(__half2 x, __half2 y)
hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h
+2
View File
@@ -38,6 +38,7 @@ extern "C"
__device__ __attribute__((const)) _Float16 __ocml_floor_f16(_Float16);
__device__ __attribute__((const))
_Float16 __ocml_fma_f16(_Float16, _Float16, _Float16);
__device__ __attribute__((const)) _Float16 __ocml_fabs_f16(_Float16);
__device__ __attribute__((const)) int __ocml_isinf_f16(_Float16);
__device__ __attribute__((const)) int __ocml_isnan_f16(_Float16);
__device__ __attribute__((pure)) _Float16 __ocml_log_f16(_Float16);
@@ -58,6 +59,7 @@ extern "C"
#endif
__device__ __attribute__((const)) __2f16 __ocml_ceil_2f16(__2f16);
__device__ __attribute__((const)) __2f16 __ocml_fabs_2f16(__2f16);
__device__ __2f16 __ocml_cos_2f16(__2f16);
__device__ __attribute__((pure)) __2f16 __ocml_exp_2f16(__2f16);
__device__ __attribute__((pure)) __2f16 __ocml_exp10_2f16(__2f16);
hipamd/include/hip/hcc_detail/hip_runtime.h
+8 -3
View File
@@ -504,9 +504,14 @@ hc_get_workitem_absolute_id(int dim)
#define __CUDA__
#include <__clang_cuda_math_forward_declares.h>
#include <__clang_cuda_complex_builtins.h>
#include <cuda_wrappers/algorithm>
#include <cuda_wrappers/complex>
#include <cuda_wrappers/new>
// Workaround for using libc++ with HIP-Clang.
// The following headers requires clang include path before standard C++ include path.
// However libc++ include path requires to be before clang include path.
// To workaround this, we pass -isystem with the parent directory of clang include
// path instead of the clang include path itself.
#include <include/cuda_wrappers/algorithm>
#include <include/cuda_wrappers/complex>
#include <include/cuda_wrappers/new>
#undef __CUDA__
#pragma pop_macro("__CUDA__")
#endif // !_OPENMP || __HIP_ENABLE_CUDA_WRAPPER_FOR_OPENMP__
hipamd/include/hip/hcc_detail/hip_runtime_api.h
+1 -1
View File
@@ -55,7 +55,7 @@ THE SOFTWARE.
#define DEPRECATED(msg) __attribute__ ((deprecated(msg)))
#endif // !defined(_MSC_VER)
#define DEPRECATED_MSG "This API is marked as deprecated and may not be supported in future releases.For more details please refer https://github.com/ROCm-Developer-Tools/HIP/tree/master/docs/markdown/hip_deprecated_api_list"
#define DEPRECATED_MSG "This API is marked as deprecated and may not be supported in future releases. For more details please refer https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_deprecated_api_list.md"
#if defined(__HCC__) && (__hcc_workweek__ < 16155)
#error("This version of HIP requires a newer version of HCC.");
hipamd/include/hip/hcc_detail/hip_vector_types.h
+93 -4
View File
@@ -34,7 +34,7 @@ THE SOFTWARE.
#include "hip/hcc_detail/host_defines.h"
#if !defined(_MSC_VER) || __clang__
#if defined(__has_attribute)
#if __has_attribute(ext_vector_type)
#define __NATIVE_VECTOR__(n, T) T __attribute__((ext_vector_type(n)))
#else
@@ -694,7 +694,7 @@ THE SOFTWARE.
typename U = T,
typename std::enable_if<std::is_signed<U>{}>::type* = nullptr>
inline __host__ __device__
HIP_vector_type operator-() noexcept
HIP_vector_type operator-() const noexcept
{
auto tmp(*this);
tmp.data = -tmp.data;
@@ -705,7 +705,7 @@ THE SOFTWARE.
typename U = T,
typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
inline __host__ __device__
HIP_vector_type operator~() noexcept
HIP_vector_type operator~() const noexcept
{
HIP_vector_type r{*this};
r.data = ~r.data;
@@ -1241,7 +1241,9 @@ DECLOP_MAKE_ONE_COMPONENT(signed long long, longlong1);
DECLOP_MAKE_TWO_COMPONENT(signed long long, longlong2);
DECLOP_MAKE_THREE_COMPONENT(signed long long, longlong3);
DECLOP_MAKE_FOUR_COMPONENT(signed long long, longlong4);
#else // defined(_MSC_VER)
#else // !defined(__has_attribute)
#if defined(_MSC_VER)
#include <mmintrin.h>
#include <xmmintrin.h>
#include <emmintrin.h>
@@ -1347,5 +1349,92 @@ typedef union { double4 data; } double3;
typedef union { __m256d data[2]; } double8;
typedef union { __m256d data[4]; } double16;
#else // !defined(_MSC_VER)
typedef union { char data; } char1;
typedef union { char data[2]; } char2;
typedef union { char data[4]; } char4;
typedef union { char data[8]; } char8;
typedef union { char data[16]; } char16;
typedef union { char4 data; } char3;
typedef union { unsigned char data; } uchar1;
typedef union { unsigned char data[2]; } uchar2;
typedef union { unsigned char data[4]; } uchar4;
typedef union { unsigned char data[8]; } uchar8;
typedef union { unsigned char data[16]; } uchar16;
typedef union { uchar4 data; } uchar3;
typedef union { short data; } short1;
typedef union { short data[2]; } short2;
typedef union { short data[4]; } short4;
typedef union { short data[8]; } short8;
typedef union { short data[16]; } short16;
typedef union { short4 data; } short3;
typedef union { unsigned short data; } ushort1;
typedef union { unsigned short data[2]; } ushort2;
typedef union { unsigned short data[4]; } ushort4;
typedef union { unsigned short data[8]; } ushort8;
typedef union { unsigned short data[16]; } ushort16;
typedef union { ushort4 data; } ushort3;
typedef union { int data; } int1;
typedef union { int data[2]; } int2;
typedef union { int data[4]; } int4;
typedef union { int data[8]; } int8;
typedef union { int data[16]; } int16;
typedef union { int4 data; } int3;
typedef union { unsigned int data; } uint1;
typedef union { unsigned int data[2]; } uint2;
typedef union { unsigned int data[4]; } uint4;
typedef union { unsigned int data[8]; } uint8;
typedef union { unsigned int data[16]; } uint16;
typedef union { uint4 data; } uint3;
typedef union { long data; } long1;
typedef union { long data[2]; } long2;
typedef union { long data[4]; } long4;
typedef union { long data[8]; } long8;
typedef union { long data[16]; } long16;
typedef union { long4 data; } long3;
typedef union { unsigned long data; } ulong1;
typedef union { unsigned long data[2]; } ulong2;
typedef union { unsigned long data[4]; } ulong4;
typedef union { unsigned long data[8]; } ulong8;
typedef union { unsigned long data[16]; } ulong16;
typedef union { ulong4 data; } ulong3;
typedef union { long long data; } longlong1;
typedef union { long long data[2]; } longlong2;
typedef union { long long data[4]; } longlong4;
typedef union { long long data[8]; } longlong8;
typedef union { long long data[16]; } longlong16;
typedef union { longlong4 data; } longlong3;
typedef union { unsigned long long data; } ulonglong1;
typedef union { unsigned long long data[2]; } ulonglong2;
typedef union { unsigned long long data[4]; } ulonglong4;
typedef union { unsigned long long data[8]; } ulonglong8;
typedef union { unsigned long long data[16]; } ulonglong16;
typedef union { ulonglong4 data; } ulonglong3;
typedef union { float data; } float1;
typedef union { float data[2]; } float2;
typedef union { float data[4]; } float4;
typedef union { float data[8]; } float8;
typedef union { float data[16]; } float16;
typedef union { float4 data; } float3;
typedef union { double data; } double1;
typedef union { double data[2]; } double2;
typedef union { double data[4]; } double4;
typedef union { double data[8]; } double8;
typedef union { double data[16]; } double16;
typedef union { double4 data; } double3;
#endif // defined(_MSC_VER)
#endif // defined(__has_attribute)
#endif
hipamd/include/hip/hcc_detail/hiprtc.h
+4
View File
@@ -28,6 +28,8 @@ extern "C" {
#include <stdlib.h>
#pragma GCC visibility push (default)
enum hiprtcResult {
HIPRTC_SUCCESS = 0,
HIPRTC_ERROR_OUT_OF_MEMORY = 1,
@@ -79,6 +81,8 @@ hiprtcResult hiprtcGetCode(hiprtcProgram prog, char* code);
hiprtcResult hiprtcGetCodeSize(hiprtcProgram prog, size_t* codeSizeRet);
#pragma GCC visibility pop
#ifdef __cplusplus
}
#endif /* __cplusplus */
hipamd/include/hip/hcc_detail/host_defines.h
+1 -1
View File
@@ -60,7 +60,7 @@ THE SOFTWARE.
*/
// _restrict is supported by the compiler
#define __shared__ tile_static
#define __constant__ __attribute__((hc))
#define __constant__ __attribute__((hc, annotate("__HIP_constant__")))
#elif defined(__clang__) && defined(__HIP__)
hipamd/include/hip/hip_runtime_api.h
-1
View File
@@ -321,7 +321,6 @@ typedef enum hipDeviceAttribute_t {
hipDeviceAttributeIntegrated, ///< iGPU
hipDeviceAttributeCooperativeLaunch, ///< Support cooperative launch
hipDeviceAttributeCooperativeMultiDeviceLaunch, ///< Support cooperative launch on multiple devices
hipDeviceAttributeMaxTexture1DWidth, ///< Maximum number of elements in 1D images
hipDeviceAttributeMaxTexture2DWidth, ///< Maximum dimension width of 2D images in image elements
hipDeviceAttributeMaxTexture2DHeight, ///< Maximum dimension height of 2D images in image elements
hipamd/include/hip/nvcc_detail/hip_runtime_api.h
+42 -1
View File
@@ -186,6 +186,7 @@ typedef struct cudaArray hipArray;
typedef struct cudaArray* hipArray_t;
typedef struct cudaArray* hipArray_const_t;
typedef struct cudaFuncAttributes hipFuncAttributes;
typedef struct cudaLaunchParams hipLaunchParams;
#define hipFunction_attribute CUfunction_attribute
#define hip_Memcpy2D CUDA_MEMCPY2D
#define hipMemcpy3DParms cudaMemcpy3DParms
@@ -860,7 +861,7 @@ inline static hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes,
}
inline hipError_t hipMemcpyWithStream(void* dst, const void* src,
inline static hipError_t hipMemcpyWithStream(void* dst, const void* src,
size_t sizeBytes, hipMemcpyKind copyKind,
hipStream_t stream) {
cudaError_t error = cudaMemcpyAsync(dst, src, sizeBytes,
@@ -1134,6 +1135,10 @@ inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t* p_prop, int dev
p_prop->integrated = cdprop.integrated;
p_prop->cooperativeLaunch = cdprop.cooperativeLaunch;
p_prop->cooperativeMultiDeviceLaunch = cdprop.cooperativeMultiDeviceLaunch;
p_prop->cooperativeMultiDeviceUnmatchedFunc = 0;
p_prop->cooperativeMultiDeviceUnmatchedGridDim = 0;
p_prop->cooperativeMultiDeviceUnmatchedBlockDim = 0;
p_prop->cooperativeMultiDeviceUnmatchedSharedMem = 0;
p_prop->maxTexture1D = cdprop.maxTexture1D;
p_prop->maxTexture2D[0] = cdprop.maxTexture2D[0];
@@ -1271,6 +1276,12 @@ inline static hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t att
case hipDeviceAttributeEccEnabled:
cdattr = cudaDevAttrEccEnabled;
break;
case hipDeviceAttributeCooperativeLaunch:
cdattr = cudaDevAttrCooperativeLaunch;
break;
case hipDeviceAttributeCooperativeMultiDeviceLaunch:
cdattr = cudaDevAttrCooperativeMultiDeviceLaunch;
break;
default:
return hipCUDAErrorTohipError(cudaErrorInvalidValue);
}
@@ -1679,6 +1690,17 @@ inline static hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_
return hipCUDAErrorTohipError(cudaGetChannelDesc(desc,array));
}
inline static hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim, dim3 blockDim,
void** kernelParams, unsigned int sharedMemBytes,
hipStream_t stream) {
return hipCUDAErrorTohipError(
cudaLaunchCooperativeKernel(f, gridDim, blockDim, kernelParams, sharedMemBytes, stream));
}
inline static hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList,
int numDevices, unsigned int flags) {
return hipCUDAErrorTohipError(cudaLaunchCooperativeKernelMultiDevice(launchParamsList, numDevices, flags));
}
#ifdef __cplusplus
}
@@ -1686,6 +1708,17 @@ inline static hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_
#ifdef __CUDACC__
template<class T>
inline static hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(int* numBlocks,
T func,
int blockSize,
size_t dynamicSMemSize) {
cudaError_t cerror;
cerror =
cudaOccupancyMaxActiveBlocksPerMultiprocessor(numBlocks, func, blockSize, dynamicSMemSize);
return hipCUDAErrorTohipError(cerror);
}
template <class T>
inline static hipError_t hipOccupancyMaxPotentialBlockSize(int* minGridSize, int* blockSize, T func,
size_t dynamicSMemSize = 0,
@@ -1742,6 +1775,14 @@ template <class T>
inline static hipChannelFormatDesc hipCreateChannelDesc() {
return cudaCreateChannelDesc<T>();
}
template <class T>
inline static hipError_t hipLaunchCooperativeKernel(T f, dim3 gridDim, dim3 blockDim,
void** kernelParams, unsigned int sharedMemBytes, hipStream_t stream) {
return hipCUDAErrorTohipError(
cudaLaunchCooperativeKernel(f, gridDim, blockDim, kernelParams, sharedMemBytes, stream));
}
#endif //__CUDACC__
#endif // HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_RUNTIME_API_H