AI/rocm-systems
2
0
Fork 0
Código Incidencias Pull Requests Acciones Paquetes Proyectos Lanzamientos Wiki Actividad

Merge 'master' into 'amd-master'

Explorar el Código 
Change-Id: Ieb80c129669e5c781fbad1e217be399915421cf2
Este commit está contenido en:
Jenkins
2018-09-17 04:09:45 -05:00
padre cc17091ca3 30b1494c3d
commit f7c25a20b6
Se han modificado 28 ficheros con 836 adiciones y 154 borrados
Descargar archivo de parche Descargar archivo de diferencias Expandir todos los archivos Contraer todos los archivos
CMakeLists.txt
+4 -2
Ver fichero
@@ -217,9 +217,11 @@ if(HIP_PLATFORM STREQUAL "hcc")
set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L/opt/rocm/profiler/CXLActivityLogger/bin/x86_64 -lCXLActivityLogger")
endif()
add_library(hip_hcc SHARED ${SOURCE_FILES_RUNTIME})
target_link_libraries(hip_hcc PRIVATE hc_am)
add_library(hip_hcc_static STATIC ${SOURCE_FILES_RUNTIME})
target_link_libraries(hip_hcc_static PRIVATE hc_am)
if(HIP_COMPILER STREQUAL "hcc")
target_link_libraries(hip_hcc PRIVATE hc_am)
target_link_libraries(hip_hcc_static PRIVATE hc_am)
endif()
add_library(hip_device STATIC ${SOURCE_FILES_DEVICE})
string(REPLACE " " ";" HCC_CXX_FLAGS_LIST ${HCC_CXX_FLAGS})
Jenkinsfile
vendido
+1 -1
Ver fichero
@@ -177,7 +177,7 @@ def docker_build_inside_image( def build_image, String inside_args, String platf
cd ${build_dir_rel}
make install -j\$(nproc)
make build_tests -i -j\$(nproc)
ctest
ctest -E hipVectorTypes
"""
// If unit tests output a junit or xunit file in the future, jenkins can parse that file
// to display test results on the dashboard
bin/hipcc
+10 -1
Ver fichero
@@ -453,6 +453,15 @@ foreach $arg (@ARGV)
system("cd $tmpdir; ar c $libBaseName $realObjs");
print $out "$tmpdir/$libBaseName\n";
}
} elsif ($line =~ m/\.o$/) {
my $fileType = `file $line`;
my $isObj = ($fileType =~ m/ELF/ or $fileType =~ m/COFF/);
if ($isObj) {
print $out "$line\n";
} else {
push (@inputs, $line);
$new_arg = "$new_arg $line";
}
} else {
print $out "$line\n";
}
@@ -725,7 +734,7 @@ if ($needLDFLAGS and not $compileOnly) {
}
$CMD .= " $toolArgs";
if ($needLDFLAGS and not $compileOnly and $HIP_PLATFORM eq "clang") {
$CMD .= " -lgcc_s -lgcc";
$CMD .= " -lgcc_s -lgcc -lpthread -lm";
}
if ($verbose & 0x1) {
docs/markdown/hip_porting_driver_api.md
+9
Ver fichero
@@ -102,6 +102,15 @@ hip-clang links device code from different translation units together. For each
hip-clang generates initializatiion and termination functions for each translation unit for host code compilation. The initialization functions call `__hipRegisterFatBinary` to register the fatbinary embeded in the ELF file. They also call `__hipRegisterFunction` and `__hipRegisterVar` to register kernel functions and device side global variables. The termination functions call `__hipUnregisterFatBinary`.
hip-clang emits a global variable `__hip_gpubin_handle` of void** type with linkonce linkage and inital value 0 for each host translation unit. Each initialization function checks `__hip_gpubin_handle` and register the fatbinary only if `__hip_gpubin_handle` is 0 and saves the return value of `__hip_gpubin_handle` to `__hip_gpubin_handle`. This is to guarantee that the fatbinary is only registered once. Similar check is done in the termination functions.
#### Kernel Launching
hip-clang supports kernel launching by CUDA `<<<>>>` syntax, hipLaunchKernel, and hipLaunchKernelGGL. The latter two are macros which expand to CUDA `<<<>>>` syntax.
In host code, hip-clang emits a stub function with the same name and arguments as the kernel. In the body of this function, hipSetupArgument is called for each kernel argument, then hipLaunchByPtr is called with a function pointer to the stub function.
When the executable or shared library is loaded by the dynamic linker, the initilization functions are called. In the initialization functions, when `__hipRegisterFatBinary` is called, the code objects containing all kernels are loaded; when `__hipRegisterFunction` is called, the stub functions are associated with the corresponding kernels in code objects.
In the host code, for the `<<<>>>` statement, hip-clang first emits call of hipConfigureCall to set up the threads and grids, then emits call of the stub function with the given arguments. In the stub function, when the runtime host API function hipLaunchByPtr is called, the real kernel associated with the stub function is launched.
### NVCC Implementation Notes
#### Interoperation between HIP and CUDA Driver
docs/markdown/hip_programming_guide.md
+15 -1
Ver fichero
@@ -91,8 +91,22 @@ Setting HCC_UNPINNED_COPY_MODE = 3, forces all unpinned transfer to use direct m
Following environment variables can be used to control the transfer thresholds:
- HCC_H2D_STAGING_THRESHOLD - Threshold in KB for H2D copy. For sizes smaller than threshold direct copy logic would be used else staging buffers logic. By default it is set to 64.
- HCC_H2D_STAGING_THRESHOLD - Threshold in KB for H2D copy. For sizes smaller than threshold direct copy logic would be used else staging buffers logic. By default it is set to 64.
- HCC_H2D_PININPLACE_THRESHOLD - Threshold in KB for H2D copy. For sizes smaller than threshold staging buffers logic would be used else PinInPlace logic. By default it is set to 4096.
- HCC_D2H_PININPLACE_THRESHOLD - Threshold in KB for D2H copy. For sizes smaller than threshold staging buffer logic would be used else PinInPlace logic. By default it is set to 1024.
## Device-Side Malloc
hip-hcc and hip-clang supports device-side malloc and free. Users can allocate
memory dynamically in a kernel. The allocated memory are in global address
space, however, different threads get different memory allocations for the same
call of malloc. The allocated memory can be accessed or freed by other threads
or other kernels. It persists in the life time of the HIP program until it is
freed.
The memory are allocated in pages. Users can define macro
`__HIP_SIZE_OF_PAGE` for controlling the page size in bytes and macro
`__HIP_NUM_PAGES` for controlling the total number of pages that can be
allocated.
hip-config.cmake.in
+3 -1
Ver fichero
@@ -48,7 +48,9 @@ set_and_check( hip_BIN_INSTALL_DIR "@PACKAGE_BIN_INSTALL_DIR@" )
set_and_check(hip_HIPCC_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipcc")
set_and_check(hip_HIPCONFIG_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipconfig")
find_dependency(hcc)
if(HIP_COMPILER STREQUAL "hcc")
find_dependency(hcc)
endif()
include( "${CMAKE_CURRENT_LIST_DIR}/hip-targets.cmake" )
set( hip_LIBRARIES hip::host hip::device)
include/hip/hcc_detail/device_functions.h
+2 -1
Ver fichero
@@ -593,7 +593,7 @@ __device__ static inline double __longlong_as_double(long long int x) {
double tmp;
__builtin_memcpy(&tmp, &x, sizeof(tmp));
return x;
return tmp;
}
__device__ static inline double __uint2double_rn(int x) { return (double)x; }
@@ -1036,4 +1036,5 @@ static inline __device__ void* memset(void* ptr, int val, size_t size) {
unsigned char val8 = static_cast<unsigned char>(val);
return __hip_hc_memset(ptr, val8, size);
}
#endif
include/hip/hcc_detail/functional_grid_launch.hpp
+44 -20
Ver fichero
@@ -51,30 +51,53 @@ inline T round_up_to_next_multiple_nonnegative(T x, T y) {
return tmp - tmp % y;
}
inline std::vector<std::uint8_t> make_kernarg() { return {}; }
inline std::vector<std::uint8_t> make_kernarg(std::vector<std::uint8_t> kernarg) { return kernarg; }
template <typename T>
inline std::vector<std::uint8_t> make_kernarg(std::vector<uint8_t> kernarg, T x) {
kernarg.resize(round_up_to_next_multiple_nonnegative(kernarg.size(), alignof(T)) + sizeof(T));
new (kernarg.data() + kernarg.size() - sizeof(T)) T{std::move(x)};
template <
std::size_t n,
typename... Ts,
typename std::enable_if<n == sizeof...(Ts)>::type* = nullptr>
inline std::vector<std::uint8_t> make_kernarg(
std::vector<std::uint8_t> kernarg, const std::tuple<Ts...>&) {
return kernarg;
}
template <typename T, typename... Ts>
inline std::vector<std::uint8_t> make_kernarg(std::vector<std::uint8_t> kernarg, T x, Ts... xs) {
return make_kernarg(make_kernarg(std::move(kernarg), std::move(x)), std::move(xs)...);
template <
std::size_t n,
typename... Ts,
typename std::enable_if<n != sizeof...(Ts)>::type* = nullptr>
inline std::vector<std::uint8_t> make_kernarg(
std::vector<std::uint8_t> kernarg, const std::tuple<Ts...>& formals) {
using T = typename std::tuple_element<n, std::tuple<Ts...>>::type;
static_assert(
!std::is_reference<T>{},
"A __global__ function cannot have a reference as one of its "
"arguments.");
#if defined(HIP_STRICT)
static_assert(
std::is_trivially_copyable<T>{},
"Only TriviallyCopyable types can be arguments to a __global__ "
"function");
#endif
kernarg.resize(round_up_to_next_multiple_nonnegative(
kernarg.size(), alignof(T)) + sizeof(T));
new (kernarg.data() + kernarg.size() - sizeof(T)) T{std::get<n>(formals)};
return make_kernarg<n + 1>(std::move(kernarg), formals);
}
template <typename... Ts>
inline std::vector<std::uint8_t> make_kernarg(Ts... xs) {
std::vector<std::uint8_t> kernarg;
kernarg.reserve(sizeof(std::tuple<Ts...>));
template <typename... Formals, typename... Actuals>
inline std::vector<std::uint8_t> make_kernarg(
void (*)(Formals...), std::tuple<Actuals...> actuals) {
static_assert(sizeof...(Formals) == sizeof...(Actuals),
"The count of formal arguments must match the count of actuals.");
return make_kernarg(std::move(kernarg), std::move(xs)...);
std::tuple<Formals...> to_formals{std::move(actuals)};
std::vector<std::uint8_t> kernarg;
kernarg.reserve(sizeof(to_formals));
return make_kernarg<0>(std::move(kernarg), to_formals);
}
void hipLaunchKernelGGLImpl(std::uintptr_t function_address, const dim3& numBlocks,
@@ -85,7 +108,8 @@ void hipLaunchKernelGGLImpl(std::uintptr_t function_address, const dim3& numBloc
template <typename... Args, typename F = void (*)(Args...)>
inline void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
std::uint32_t sharedMemBytes, hipStream_t stream, Args... args) {
auto kernarg = hip_impl::make_kernarg(std::move(args)...);
auto kernarg = hip_impl::make_kernarg(
kernel, std::tuple<Args...>{std::move(args)...});
std::size_t kernarg_size = kernarg.size();
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, kernarg.data(), HIP_LAUNCH_PARAM_BUFFER_SIZE,
@@ -100,4 +124,4 @@ inline void hipLaunchKernel(F kernel, const dim3& numBlocks, const dim3& dimBloc
std::uint32_t groupMemBytes, hipStream_t stream, Args... args) {
hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, groupMemBytes, stream, hipLaunchParm{},
std::move(args)...);
}
}
include/hip/hcc_detail/hip_fp16.h
+13 -13
Ver fichero
@@ -635,37 +635,37 @@ THE SOFTWARE.
// TODO: rounding behaviour is not correct.
// float -> half | half2
inline
__device__
__device__ __host__
__half __float2half(float x)
{
return __half_raw{static_cast<_Float16>(x)};
}
inline
__device__
__device__ __host__
__half __float2half_rn(float x)
{
return __half_raw{static_cast<_Float16>(x)};
}
inline
__device__
__device__ __host__
__half __float2half_rz(float x)
{
return __half_raw{static_cast<_Float16>(x)};
}
inline
__device__
__device__ __host__
__half __float2half_rd(float x)
{
return __half_raw{static_cast<_Float16>(x)};
}
inline
__device__
__device__ __host__
__half __float2half_ru(float x)
{
return __half_raw{static_cast<_Float16>(x)};
}
inline
__device__
__device__ __host__
__half2 __float2half2_rn(float x)
{
return __half2_raw{
@@ -673,14 +673,14 @@ THE SOFTWARE.
static_cast<_Float16>(x), static_cast<_Float16>(x)}};
}
inline
__device__
__device__ __host__
__half2 __floats2half2_rn(float x, float y)
{
return __half2_raw{_Float16_2{
static_cast<_Float16>(x), static_cast<_Float16>(y)}};
}
inline
__device__
__device__ __host__
__half2 __float22half2_rn(float2 x)
{
return __floats2half2_rn(x.x, x.y);
@@ -688,25 +688,25 @@ THE SOFTWARE.
// half | half2 -> float
inline
__device__
__device__ __host__
float __half2float(__half x)
{
return static_cast<__half_raw>(x).data;
}
inline
__device__
__device__ __host__
float __low2float(__half2 x)
{
return static_cast<__half2_raw>(x).data.x;
}
inline
__device__
__device__ __host__
float __high2float(__half2 x)
{
return static_cast<__half2_raw>(x).data.y;
}
inline
__device__
__device__ __host__
float2 __half22float2(__half2 x)
{
return make_float2(
@@ -1633,4 +1633,4 @@ THE SOFTWARE.
#endif // defined(__cplusplus)
#elif defined(__GNUC__)
#include "hip_fp16_gcc.h"
#endif // !defined(__clang__) && defined(__GNUC__)
#endif // !defined(__clang__) && defined(__GNUC__)
include/hip/hcc_detail/hip_memory.h
+105
Ver fichero
@@ -0,0 +1,105 @@
/*
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.
*/
#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_MEMORY_H
#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_MEMORY_H
// Implementation of malloc and free device functions.
// HIP heap is implemented as a global array with fixed size. Users may define
// __HIP_SIZE_OF_PAGE and __HIP_NUM_PAGES to have a larger heap.
// Size of page in bytes.
#ifndef __HIP_SIZE_OF_PAGE
#define __HIP_SIZE_OF_PAGE 64
#endif
// Total number of pages
#ifndef __HIP_NUM_PAGES
#define __HIP_NUM_PAGES (16 * 64 * 64)
#endif
#define __HIP_SIZE_OF_HEAP (__HIP_NUM_PAGES * __HIP_SIZE_OF_PAGE)
__attribute__((weak)) __device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP];
__attribute__((weak)) __device__
uint32_t __hip_device_page_flag[__HIP_NUM_PAGES];
extern "C" inline __device__ void* __hip_malloc(size_t size) {
char* heap = (char*)__hip_device_heap;
if (size > __HIP_SIZE_OF_HEAP) {
return (void*)nullptr;
}
uint32_t totalThreads =
hipBlockDim_x * hipGridDim_x * hipBlockDim_y
* hipGridDim_y * hipBlockDim_z * hipGridDim_z;
uint32_t currentWorkItem = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x
+ (hipThreadIdx_y + hipBlockDim_y * hipBlockIdx_y) * hipBlockDim_x
+ (hipThreadIdx_z + hipBlockDim_z * hipBlockIdx_z) * hipBlockDim_x
* hipBlockDim_y;
uint32_t numHeapsPerWorkItem = __HIP_NUM_PAGES / totalThreads;
uint32_t heapSizePerWorkItem = __HIP_SIZE_OF_HEAP / totalThreads;
uint32_t stride = size / __HIP_SIZE_OF_PAGE;
uint32_t start = numHeapsPerWorkItem * currentWorkItem;
uint32_t k = 0;
while (__hip_device_page_flag[k] > 0) {
k++;
}
for (uint32_t i = 0; i < stride - 1; i++) {
__hip_device_page_flag[i + start + k] = 1;
}
__hip_device_page_flag[start + stride - 1 + k] = 2;
void* ptr = (void*)(heap
+ heapSizePerWorkItem * currentWorkItem + k * __HIP_SIZE_OF_PAGE);
return ptr;
}
extern "C" inline __device__ void* __hip_free(void* ptr) {
if (ptr == nullptr) {
return nullptr;
}
uint32_t offsetByte = (uint64_t)ptr - (uint64_t)__hip_device_heap;
uint32_t offsetPage = offsetByte / __HIP_SIZE_OF_PAGE;
while (__hip_device_page_flag[offsetPage] != 0) {
if (__hip_device_page_flag[offsetPage] == 2) {
__hip_device_page_flag[offsetPage] = 0;
offsetPage++;
break;
} else {
__hip_device_page_flag[offsetPage] = 0;
offsetPage++;
}
}
return nullptr;
}
#endif // HIP_INCLUDE_HIP_HCC_DETAIL_HIP_MEMORY_H
include/hip/hcc_detail/hip_runtime.h
+22 -28
Ver fichero
@@ -199,16 +199,6 @@ __device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask
#endif //__HIP_ARCH_GFX803__ == 1
__device__ inline static int min(int arg1, int arg2) {
return (arg1 < arg2) ? arg1 : arg2;
}
__device__ inline static int max(int arg1, int arg2) {
return (arg1 > arg2) ? arg1 : arg2;
}
__host__ inline static int min(int arg1, int arg2) { return std::min(arg1, arg2); }
__host__ inline static int max(int arg1, int arg2) { return std::max(arg1, arg2); }
#endif // __HCC_OR_HIP_CLANG__
#if defined __HCC__
@@ -258,25 +248,22 @@ static constexpr Coordinates<hc_get_workitem_id> threadIdx;
#endif // defined __HCC__
#if __HCC_OR_HIP_CLANG__
extern "C" __device__ void* __hip_hc_malloc(size_t);
extern "C" __device__ void* __hip_hc_free(void* ptr);
extern "C" __device__ void* __hip_malloc(size_t);
extern "C" __device__ void* __hip_free(void* ptr);
static inline __device__ void* malloc(size_t size) { return __hip_hc_malloc(size); }
static inline __device__ void* free(void* ptr) { return __hip_hc_free(ptr); }
static inline __device__ void* malloc(size_t size) { return __hip_malloc(size); }
static inline __device__ void* free(void* ptr) { return __hip_free(ptr); }
#ifdef __HCC_ACCELERATOR__
#ifdef HC_FEATURE_PRINTF
#if defined(__HCC_ACCELERATOR__) && defined(HC_FEATURE_PRINTF)
template <typename... All>
static inline __device__ void printf(const char* format, All... all) {
hc::printf(format, all...);
}
#else
#elif defined(__HCC_ACCELERATOR__) || __HIP__
template <typename... All>
static inline __device__ void printf(const char* format, All... all) {}
#endif
#endif
#endif //__HCC_OR_HIP_CLANG__
#ifdef __HCC__
@@ -346,15 +333,18 @@ extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, gri
typedef int hipLaunchParm;
#define hipLaunchKernel(kernelName, numblocks, numthreads, memperblock, streamId, ...) \
do { \
kernelName<<<numblocks, numthreads, memperblock, streamId>>>(0, ##__VA_ARGS__); \
} while (0)
template <typename... Args, typename F = void (*)(Args...)>
inline void hipLaunchKernelGGL(F kernelName, const dim3& numblocks, const dim3& numthreads,
unsigned memperblock, hipStream_t streamId, Args... args) {
kernelName<<<numblocks, numthreads, memperblock, streamId>>>(args...);
}
#define hipLaunchKernelGGL(kernelName, numblocks, numthreads, memperblock, streamId, ...) \
do { \
kernelName<<<numblocks, numthreads, memperblock, streamId>>>(__VA_ARGS__); \
} while (0)
template <typename... Args, typename F = void (*)(hipLaunchParm, Args...)>
inline void hipLaunchKernel(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
std::uint32_t groupMemBytes, hipStream_t stream, Args... args) {
hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, groupMemBytes, stream, hipLaunchParm{},
std::move(args)...);
}
#include <hip/hip_runtime_api.h>
@@ -436,6 +426,8 @@ extern const __device__ __attribute__((weak)) __hip_builtin_gridDim_t gridDim;
#define hipGridDim_y gridDim.y
#define hipGridDim_z gridDim.z
#include <hip/hcc_detail/math_functions.h>
#if __HIP_HCC_COMPAT_MODE__
// Define HCC work item functions in terms of HIP builtin variables.
#pragma push_macro("__DEFINE_HCC_FUNC")
@@ -471,10 +463,10 @@ hc_get_workitem_absolute_id(int dim)
#include <__clang_cuda_complex_builtins.h>
#include <cuda_wrappers/algorithm>
#include <cuda_wrappers/complex>
#include <cuda_wrappers/new>
#undef __CUDA__
#pragma pop_macro("__CUDA__")
#include <hip/hcc_detail/math_functions.h>
hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
@@ -486,4 +478,6 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
#endif // defined(__clang__) && defined(__HIP__)
#include <hip/hcc_detail/hip_memory.h>
#endif // HIP_HCC_DETAIL_RUNTIME_H
include/hip/hcc_detail/math_functions.h
+71
Ver fichero
@@ -1304,6 +1304,66 @@ float func(float x, int y) \
}
__DEF_FLOAT_FUN2I(scalbn)
#if __HCC__
template<class T>
__DEVICE__ inline static T min(T arg1, T arg2) {
return (arg1 < arg2) ? arg1 : arg2;
}
__DEVICE__ inline static uint32_t min(uint32_t arg1, int32_t arg2) {
return min(arg1, (uint32_t) arg2);
}
/*__DEVICE__ inline static uint32_t min(int32_t arg1, uint32_t arg2) {
return min((uint32_t) arg1, arg2);
}
__DEVICE__ inline static uint64_t min(uint64_t arg1, int64_t arg2) {
return min(arg1, (uint64_t) arg2);
}
__DEVICE__ inline static uint64_t min(int64_t arg1, uint64_t arg2) {
return min((uint64_t) arg1, arg2);
}
__DEVICE__ inline static unsigned long long min(unsigned long long arg1, long long arg2) {
return min(arg1, (unsigned long long) arg2);
}
__DEVICE__ inline static unsigned long long min(long long arg1, unsigned long long arg2) {
return min((unsigned long long) arg1, arg2);
}*/
template<class T>
__DEVICE__ inline static T max(T arg1, T arg2) {
return (arg1 > arg2) ? arg1 : arg2;
}
__DEVICE__ inline static uint32_t max(uint32_t arg1, int32_t arg2) {
return max(arg1, (uint32_t) arg2);
}
__DEVICE__ inline static uint32_t max(int32_t arg1, uint32_t arg2) {
return max((uint32_t) arg1, arg2);
}
/*__DEVICE__ inline static uint64_t max(uint64_t arg1, int64_t arg2) {
return max(arg1, (uint64_t) arg2);
}
__DEVICE__ inline static uint64_t max(int64_t arg1, uint64_t arg2) {
return max((uint64_t) arg1, arg2);
}
__DEVICE__ inline static unsigned long long max(unsigned long long arg1, long long arg2) {
return max(arg1, (unsigned long long) arg2);
}
__DEVICE__ inline static unsigned long long max(long long arg1, unsigned long long arg2) {
return max((unsigned long long) arg1, arg2);
}*/
#else
__DEVICE__ inline static int min(int arg1, int arg2) {
return (arg1 < arg2) ? arg1 : arg2;
}
__DEVICE__ inline static int max(int arg1, int arg2) {
return (arg1 > arg2) ? arg1 : arg2;
}
__DEVICE__
inline
float max(float x, float y) {
@@ -1331,6 +1391,17 @@ double min(double x, double y) {
__HIP_OVERLOAD2(double, max)
__HIP_OVERLOAD2(double, min)
#endif
__host__ inline static int min(int arg1, int arg2) {
return std::min(arg1, arg2);
}
__host__ inline static int max(int arg1, int arg2) {
return std::max(arg1, arg2);
}
#pragma pop_macro("__DEF_FLOAT_FUN")
#pragma pop_macro("__DEF_FLOAT_FUN2")
#pragma pop_macro("__DEF_FLOAT_FUN2I")
packaging/hip-targets-release.cmake
+14
Ver fichero
@@ -7,22 +7,36 @@ set(CMAKE_IMPORT_FILE_VERSION 1)
# Import target "hip::hip_hcc_static" for configuration "Release"
set_property(TARGET hip::hip_hcc_static APPEND PROPERTY IMPORTED_CONFIGURATIONS RELEASE)
if(HIP_COMPILER STREQUAL "clang")
set_target_properties(hip::hip_hcc_static PROPERTIES
IMPORTED_LINK_INTERFACE_LANGUAGES_RELEASE "CXX"
IMPORTED_LOCATION_RELEASE "/opt/rocm/hip/lib/libhip_hcc_static.a"
)
else()
set_target_properties(hip::hip_hcc_static PROPERTIES
IMPORTED_LINK_INTERFACE_LANGUAGES_RELEASE "CXX"
IMPORTED_LINK_INTERFACE_LIBRARIES_RELEASE "hc_am"
IMPORTED_LOCATION_RELEASE "/opt/rocm/hip/lib/libhip_hcc_static.a"
)
endif()
list(APPEND _IMPORT_CHECK_TARGETS hip::hip_hcc_static )
list(APPEND _IMPORT_CHECK_FILES_FOR_hip::hip_hcc_static "/opt/rocm/hip/lib/libhip_hcc_static.a" )
# Import target "hip::hip_hcc" for configuration "Release"
set_property(TARGET hip::hip_hcc APPEND PROPERTY IMPORTED_CONFIGURATIONS RELEASE)
if(HIP_COMPILER STREQUAL "clang")
set_target_properties(hip::hip_hcc PROPERTIES
IMPORTED_LOCATION_RELEASE "/opt/rocm/hip/lib/libhip_hcc.so"
IMPORTED_SONAME_RELEASE "libhip_hcc.so"
)
else()
set_target_properties(hip::hip_hcc PROPERTIES
IMPORTED_LINK_INTERFACE_LIBRARIES_RELEASE "hcc::hccrt;hcc::hc_am"
IMPORTED_LOCATION_RELEASE "/opt/rocm/hip/lib/libhip_hcc.so"
IMPORTED_SONAME_RELEASE "libhip_hcc.so"
)
endif()
list(APPEND _IMPORT_CHECK_TARGETS hip::hip_hcc )
list(APPEND _IMPORT_CHECK_FILES_FOR_hip::hip_hcc "/opt/rocm/hip/lib/libhip_hcc.so" )
packaging/hip-targets.cmake
+6
Ver fichero
@@ -75,9 +75,15 @@ set_target_properties(hip::host PROPERTIES
# Create imported target hip::device
add_library(hip::device INTERFACE IMPORTED)
if(HIP_COMPILER STREQUAL "clang")
set_target_properties(hip::device PROPERTIES
INTERFACE_LINK_LIBRARIES "hip::host;hip::hip_device"
)
else()
set_target_properties(hip::device PROPERTIES
INTERFACE_LINK_LIBRARIES "hip::host;hip::hip_device;hcc::hccrt;hcc::hc_am"
)
endif()
if(CMAKE_VERSION VERSION_LESS 3.0.0)
message(FATAL_ERROR "This file relies on consumers using CMake 3.0.0 or greater.")
samples/0_Intro/module_api/launchKernelHcc.cpp
+1 -1
Ver fichero
@@ -34,7 +34,7 @@ THE SOFTWARE.
#define LEN 64
#define SIZE LEN << 2
#define fileName "vcpy_kernel.code.adipose"
#define fileName "vcpy_kernel.code"
#define kernel_name "hello_world"
#define HIP_CHECK(status) \
src/device_util.cpp
-64
Ver fichero
@@ -28,70 +28,6 @@ THE SOFTWARE.
#include "hip/hip_runtime.h"
#include <atomic>
//=================================================================================================
/*
Implementation of malloc and free device functions.
This is the best place to put them because the device
global variables need to be initialized at the start.
*/
__device__ char gpuHeap[SIZE_OF_HEAP];
__device__ uint32_t gpuFlags[NUM_PAGES];
__device__ void* __hip_hc_malloc(size_t size) {
char* heap = (char*)gpuHeap;
if (size > SIZE_OF_HEAP) {
return (void*)nullptr;
}
uint32_t totalThreads =
blockDim.x * gridDim.x * blockDim.y * gridDim.y * blockDim.z * gridDim.z;
uint32_t currentWorkItem = threadIdx.x + blockDim.x * blockIdx.x;
uint32_t numHeapsPerWorkItem = NUM_PAGES / totalThreads;
uint32_t heapSizePerWorkItem = SIZE_OF_HEAP / totalThreads;
uint32_t stride = size / SIZE_OF_PAGE;
uint32_t start = numHeapsPerWorkItem * currentWorkItem;
uint32_t k = 0;
while (gpuFlags[k] > 0) {
k++;
}
for (uint32_t i = 0; i < stride - 1; i++) {
gpuFlags[i + start + k] = 1;
}
gpuFlags[start + stride - 1 + k] = 2;
void* ptr = (void*)(heap + heapSizePerWorkItem * currentWorkItem + k * SIZE_OF_PAGE);
return ptr;
}
__device__ void* __hip_hc_free(void* ptr) {
if (ptr == nullptr) {
return nullptr;
}
uint32_t offsetByte = (uint64_t)ptr - (uint64_t)gpuHeap;
uint32_t offsetPage = offsetByte / SIZE_OF_PAGE;
while (gpuFlags[offsetPage] != 0) {
if (gpuFlags[offsetPage] == 2) {
gpuFlags[offsetPage] = 0;
offsetPage++;
break;
} else {
gpuFlags[offsetPage] = 0;
offsetPage++;
}
}
return nullptr;
}
// abort
__device__ void abort() { return hc::abort(); }
src/device_util.h
-11
Ver fichero
@@ -29,14 +29,6 @@ THE SOFTWARE.
Heap size computation for malloc and free device functions.
*/
#define NUM_PAGES_PER_THREAD 16
#define SIZE_OF_PAGE 64
#define NUM_THREADS_PER_CU 64
#define NUM_CUS_PER_GPU 64 // Specific for r9 Nano
#define NUM_PAGES NUM_PAGES_PER_THREAD* NUM_THREADS_PER_CU* NUM_CUS_PER_GPU
#define SIZE_MALLOC NUM_PAGES* SIZE_OF_PAGE
#define SIZE_OF_HEAP SIZE_MALLOC
#define HIP_SQRT_2 1.41421356237
#define HIP_SQRT_PI 1.77245385091
@@ -62,9 +54,6 @@ THE SOFTWARE.
#define HIP_PI 3.14159265358979323846
__device__ void* __hip_hc_malloc(size_t size);
__device__ void* __hip_hc_free(void* ptr);
__device__ float __hip_erfinvf(float x);
__device__ double __hip_erfinv(double x);
src/hip_clang.cpp
+2
Ver fichero
@@ -58,6 +58,7 @@ __hipRegisterFatBinary(const void* data)
{
HIP_INIT();
tprintf(DB_FB, "Enter __hipRegisterFatBinary(%p)\n", data);
const __CudaFatBinaryWrapper* fbwrapper = reinterpret_cast<const __CudaFatBinaryWrapper*>(data);
if (fbwrapper->magic != __hipFatMAGIC2 || fbwrapper->version != 1) {
return nullptr;
@@ -113,6 +114,7 @@ __hipRegisterFatBinary(const void* data)
}
}
tprintf(DB_FB, "__hipRegisterFatBinary succeeds and returns %p\n", modules);
return modules;
}
src/hip_device.cpp
+1 -1
Ver fichero
@@ -99,7 +99,7 @@ hipError_t hipDeviceGetLimit(size_t* pValue, hipLimit_t limit) {
return ihipLogStatus(hipErrorInvalidValue);
}
if (limit == hipLimitMallocHeapSize) {
*pValue = (size_t)SIZE_OF_HEAP;
*pValue = (size_t)__HIP_SIZE_OF_HEAP;
return ihipLogStatus(hipSuccess);
} else {
return ihipLogStatus(hipErrorUnsupportedLimit);
src/hip_hcc.cpp
+1 -1
Ver fichero
@@ -1228,7 +1228,7 @@ void HipReadEnv() {
READ_ENV_C(release, HIP_DB, 0,
"Print debug info. Bitmask (HIP_DB=0xff) or flags separated by '+' "
"(HIP_DB=api+sync+mem+copy)",
"(HIP_DB=api+sync+mem+copy+fatbin)",
HIP_DB_callback);
if ((HIP_DB & (1 << DB_API)) && (HIP_TRACE_API == 0)) {
// Set HIP_TRACE_API default before we read it, so it is printed correctly.
src/hip_hcc_internal.h
+3 -1
Ver fichero
@@ -223,7 +223,8 @@ extern const char* API_COLOR_END;
#define DB_MEM 2 /* 0x04 - trace memory allocation / deallocation */
#define DB_COPY 3 /* 0x08 - trace memory copy and peer commands. . */
#define DB_WARN 4 /* 0x10 - warn about sub-optimal or shady behavior */
#define DB_MAX_FLAG 5
#define DB_FB 5 /* 0x20 - trace loading fat binary */
#define DB_MAX_FLAG 6
// When adding a new debug flag, also add to the char name table below.
//
//
@@ -237,6 +238,7 @@ struct DbName {
static const DbName dbName[] = {
{KGRN, "api"}, // not used,
{KYEL, "sync"}, {KCYN, "mem"}, {KMAG, "copy"}, {KRED, "warn"},
{KBLU, "fatbin"},
};
tests/src/compiler/hipClassKernel.cpp
+87 -2
Ver fichero
@@ -26,6 +26,66 @@ THE SOFTWARE.
*/
#include "hipClassKernel.h"
#ifdef ENABLE_OVERLOAD_OVERRIDE_TESTS
__global__ void
ovrdClassKernel(bool* result_ecd){
int tid = threadIdx.x + blockIdx.x * blockDim.x;
testOvrD tobj1;
result_ecd[tid] = (tobj1.ovrdFunc1() == 30);
}
void HipClassTests::TestForOverride(void){
bool *result_ecd, *result_ech;
result_ech = HipClassTests::AllocateHostMemory();
result_ecd = HipClassTests::AllocateDeviceMemory();
hipLaunchKernelGGL(ovrdClassKernel,
dim3(BLOCKS),
dim3(THREADS_PER_BLOCK),
0,
0,
result_ecd);
HipClassTests::VerifyResult(result_ech,result_ecd);
HipClassTests::FreeMem(result_ech,result_ecd);
}
__global__ void
ovldClassKernel(bool* result_ecd){
int tid = threadIdx.x + blockIdx.x * blockDim.x;
testFuncOvld tfo1;
result_ecd[tid] = (tfo1.func1(10) == 20)
&& (tfo1.func1(10,10) == 30);
}
void HipClassTests::TestForOverload(void){
bool *result_ecd, *result_ech;
result_ech = HipClassTests::AllocateHostMemory();
result_ecd = HipClassTests::AllocateDeviceMemory();
hipLaunchKernelGGL(ovldClassKernel,
dim3(BLOCKS),
dim3(THREADS_PER_BLOCK),
0,
0,
result_ecd);
HipClassTests::VerifyResult(result_ech,result_ecd);
HipClassTests::FreeMem(result_ech,result_ecd);
}
#endif
#ifdef ENABLE_FRIEND_TEST
// check for friend
__global__ void
friendClassKernel(bool* result_ecd){
int tid = threadIdx.x + blockIdx.x * blockDim.x;
testFrndB tfb1;
result_ecd[tid] = (tfb1.showA() == 10);
}
#endif
// check sizeof empty class is 1
__global__ void
emptyClassKernel(bool* result_ecd) {
@@ -209,6 +269,20 @@ void HipClassTests::TestForConsrtDesrt(){
}
#endif
#ifdef ENABLE_FRIEND_TEST
void HipClassTests::TestForFriend(void){
bool *result_ecd, *result_ech;
result_ech = HipClassTests::AllocateHostMemory();
result_ecd = HipClassTests::AllocateDeviceMemory();
hipLaunchKernelGGL(friendClassKernel,
dim3(BLOCKS),
dim3(THREADS_PER_BLOCK),
0,
0,
result_ecd);
}
#endif
bool* HipClassTests::AllocateHostMemory(void){
bool *result_ech;
HIPCHECK(hipHostMalloc(&result_ech,
@@ -253,6 +327,19 @@ int main(){
test_passed(TestForClassSize);
classTests.TestForPassByValue();
test_passed(TestForPassByValue);
#ifdef ENABLE_OVERLOAD_OVERRIDE_TESTS
classTests.TestForOverload();
test_passed(TestForOverload);
classTests.TestForOverride();
test_passed(TestForOverride);
#endif
#ifdef ENABLE_FRIEND_TEST
classTests.TestForFriend();
test_passed(TestForFriend);
#endif
// classTests.TestForMallocPassByValue();
// test_passed(TestForMallocPassByValue); #this test is crashing
@@ -261,8 +348,6 @@ int main(){
test_passed(TestForVirtualClassSize);
#endif
#ifdef ENABLE_DESTRUCTOR_TEST
classTests.TestForConsrtDesrt();
test_passed(TestForConsrtDesrt);
tests/src/compiler/hipClassKernel.h
+55 -1
Ver fichero
@@ -31,10 +31,62 @@ static const int BLOCKS = 512;
static const int THREADS_PER_BLOCK = 1;
static const int ENABLE_DESTRUCTOR_TEST = 0;
static const int ENABLE_VIRTUAL_TESTS = 0;
static const int ENABLE_FRIEND_TEST = 0;
static const int ENABLE_OVERLAD_OVERRIDE_TESTS = 0;
size_t NBOOL = BLOCKS * sizeof(bool);
#define test_passed(test_name) printf("%s %s PASSED!%s\n", KGRN, #test_name, KNRM);
#ifdef ENABLE_OVERLOAD_OVERRIDE_TESTS
class testFuncOvld{
public:
int __host__ __device__ func1(int a){
return a + 10;
}
int __host__ __device__ func1(int a , int b){
return a + b + 10;
}
};
class testOvrB{
public:
int __host__ __device__ ovrdFunc1(){
return 10;
}
};
class testOvrD: public testOvrB{
public:
int __host__ __device__ ovrdFunc1(){
int x = testOvrB::ovrdFunc1();
return x + 20;
}
};
#endif
#ifdef ENABLE_FRIEND_TEST
class testFrndA{
private:
int fa = 10;
public:
friend class testFrndB;
};
class testFrndB{
public:
__host__ __device__ int showA(){
testFrndA x;
return x.fa;
}
};
#endif
class testClassEmpty {};
@@ -177,7 +229,9 @@ class HipClassTests{
void TestForPassByValue(void);
void TestForMallocPassByValue(void);
void TestForConsrtDesrt(void);
void TestForOverload(void);
void TestForOverride(void);
bool* AllocateHostMemory(void);
bool* AllocateDeviceMemory(void);
void VerifyResult(bool* result_ech, bool* result_ecd);
tests/src/deviceLib/hipDeviceMalloc.cpp
+188
Ver fichero
@@ -0,0 +1,188 @@
/*
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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s NVCC_OPTIONS -std=c++11
* RUN: %t EXCLUDE_HIP_PLATFORM nvcc
* HIT_END
*/
#include "test_common.h"
#include <iostream>
#include <complex>
// Tolerance for error
const double tolerance = 1e-6;
const bool verbose = false;
#define BLKDIM_X 64
#define BLKDIM_Y 1
#define BLKDIM_Z 1
#define NUM_BLK_X 1
#define NUM_BLK_Y 1
#define NUM_BLK_Z 1
#define LEN (BLKDIM_X * BLKDIM_Y * BLKDIM_Z * NUM_BLK_X * NUM_BLK_Y * NUM_BLK_Z)
#define ALL_FUN \
OP(add) \
OP(sub) \
OP(mul) \
OP(div)
#define OP(x) CK_##x,
enum CalcKind {
ALL_FUN
};
#undef OP
#define OP(x) case CK_##x: return #x;
std::string getName(enum CalcKind CK) {
switch(CK){
ALL_FUN
}
}
#undef OP
// Calculates function.
// If the function has one argument, B is ignored.
#define ONE_ARG(func) \
case CK_##func: \
return std::func(A);
template <typename FloatT>
__device__ __host__ FloatT calc(FloatT A, FloatT B, enum CalcKind CK) {
switch (CK) {
case CK_add:
return A + B;
case CK_sub:
return A - B;
case CK_mul:
return A * B;
case CK_div:
return A / B;
}
}
// Allocate memory in kernel and save the address to pA and pB.
// Copy value from A, B to allocated memory.
template <typename FloatT>
__global__ void kernel_alloc(FloatT* A, FloatT* B, FloatT** pA, FloatT** pB) {
int tx = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x
+ (hipThreadIdx_y + hipBlockDim_y * hipBlockIdx_y) * hipBlockDim_x
+ (hipThreadIdx_z + hipBlockDim_z * hipBlockIdx_z) * hipBlockDim_x
* hipBlockDim_y;
if (tx == 0) {
*pA = (FloatT*)malloc(sizeof(FloatT) * LEN);
*pB = (FloatT*)malloc(sizeof(FloatT) * LEN);
for (int i = 0; i < LEN; i++) {
(*pA)[i] = A[i];
(*pB)[i] = B[i];
}
}
}
// Do calculation using values saved in allocated memmory. pA, pB are buffers
// containing the address of the device-side allocated array.
template <typename FloatT>
__global__ void kernel_free(FloatT** pA, FloatT** pB, FloatT* C, enum CalcKind CK) {
int tx = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x
+ (hipThreadIdx_y + hipBlockDim_y * hipBlockIdx_y) * hipBlockDim_x
+ (hipThreadIdx_z + hipBlockDim_z * hipBlockIdx_z) * hipBlockDim_x
* hipBlockDim_y;
C[tx] = calc<FloatT>((*pA)[tx], (*pB)[tx], CK);
if (tx == 0) {
free(*pA);
free(*pB);
}
}
template<typename FloatT>
void test() {
FloatT *A, *Ad, *B, *Bd, *C, *Cd, *D;
A = new FloatT[LEN];
B = new FloatT[LEN];
C = new FloatT[LEN];
D = new FloatT[LEN];
hipMalloc((void**)&Ad, sizeof(FloatT) * LEN);
hipMalloc((void**)&Bd, sizeof(FloatT) * LEN);
hipMalloc((void**)&Cd, sizeof(FloatT) * LEN);
for (uint32_t i = 0; i < LEN; i++) {
A[i] = (i + 1) * 1.0f;
B[i] = A[i];
C[i] = A[i];
}
hipMemcpy(Ad, A, sizeof(FloatT) * LEN, hipMemcpyHostToDevice);
hipMemcpy(Bd, B, sizeof(FloatT) * LEN, hipMemcpyHostToDevice);
// Run kernel for a calculation kind and verify by comparing with host
// calculation result. Returns false if fails.
auto test_fun = [&](enum CalcKind CK) {
// kernel_alloc allocates memory on device side and initialize it.
// kernel_free uses allocated memory from kernel_alloc and does the
// calculation then free the memory.
// pA and pB are buffers to pass the device-side allocated memory address
// from kernel_alloc to kernel_free.
FloatT **pA, **pB;
hipMalloc((FloatT***)&pA, sizeof(FloatT*));
hipMalloc((FloatT***)&pB, sizeof(FloatT*));
dim3 blkDim(BLKDIM_X, BLKDIM_Y, BLKDIM_Z);
dim3 numBlk(NUM_BLK_X, NUM_BLK_Y, NUM_BLK_Z);
hipLaunchKernelGGL(kernel_alloc<FloatT>, numBlk, blkDim, 0, 0,
Ad, Bd, pA, pB);
hipDeviceSynchronize();
hipLaunchKernelGGL(kernel_free<FloatT>, numBlk, blkDim, 0, 0,
pA, pB, Cd, CK);
hipMemcpy(C, Cd, sizeof(FloatT) * LEN, hipMemcpyDeviceToHost);
hipFree(pA);
hipFree(pB);
for (int i = 0; i < LEN; i++) {
FloatT Expected = calc(A[i], B[i], CK);
FloatT error = std::abs(C[i] - Expected);
if (std::abs(Expected) > tolerance) error /= std::abs(Expected);
bool pass = error < tolerance;
if (verbose || !pass) {
std::cout << "Function: " << getName(CK) << " Operands: " << A[i] << " " << B[i]
<< " Result: " << C[i] << " Expected: " << Expected << " Error: " << error
<< " Pass: " << pass << std::endl;
}
if (!pass)
return false;
}
return true;
};
#define OP(x) assert(test_fun(CK_##x));
ALL_FUN
#undef OP
hipFree(Ad);
hipFree(Bd);
hipFree(Cd);
delete[] A;
delete[] B;
delete[] C;
delete[] D;
}
int main() {
test<float>();
test<double>();
passed();
return 0;
}
tests/src/deviceLib/hipTestNew.cpp
+71
Ver fichero
@@ -0,0 +1,71 @@
/*
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.
*/
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* HIT_END
*/
#include <hip/hip_runtime.h>
#include <hip/hip_runtime_api.h>
#include <iostream>
#include "test_common.h"
#define HIP_ASSERT(status) assert(status == hipSuccess)
#define LEN 512
#define SIZE 2048
struct TestPlacementNew {
class A {
public:
__device__ A() {
a = threadIdx.x + blockIdx.x * blockDim.x;
}
private:
int a;
};
static __global__ void kernel(int* Ad) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
new(Ad+tid) A();
}
void run() {
int *A, *Ad;
A = new int[LEN];
for (unsigned i = 0; i < LEN; i++) {
A[i] = 0;
}
HIP_ASSERT(hipMalloc((void**)&Ad, SIZE));
hipLaunchKernelGGL(kernel, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0, Ad);
HIP_ASSERT(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost));
for (unsigned i = 0; i < LEN; i++) {
assert(i == A[i]);
}
}
};
int main() {
TestPlacementNew().run();
passed();
}
tests/src/hipEnvVarDriver.cpp
+9 -2
Ver fichero
@@ -38,7 +38,10 @@ int getDeviceNumber() {
string str;
std::this_thread::sleep_for(std::chrono::milliseconds(10));
if (!(in = popen("./directed_tests/hipEnvVar -c", "r"))) {
return 1;
// Check at same level
if (!(in = popen("./hipEnvVar -c", "r"))) {
return 1;
}
}
while (fgets(buff, 512, in) != NULL) {
cout << buff;
@@ -54,7 +57,11 @@ void getDevicePCIBusNumRemote(int deviceID, char* pciBusID) {
str += std::to_string(deviceID);
std::this_thread::sleep_for(std::chrono::milliseconds(10));
if (!(in = popen(str.c_str(), "r"))) {
exit(1);
// Check at same level
if (!(in = popen("./hipEnvVar -d ", "r"))) {
exit(1);
}
}
while (fgets(pciBusID, 100, in) != NULL) {
cout << pciBusID;
tests/src/kernel/hipTestGlobalVariable.cpp
+97
Ver fichero
@@ -0,0 +1,97 @@
/*
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.
*/
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* RUN: %t
* HIT_END
*/
#include <hip/hip_runtime.h>
#include <hip/hip_runtime_api.h>
#include <iostream>
#include "test_common.h"
#define HIP_ASSERT(status) assert(status == hipSuccess)
#define LEN 512
#define SIZE 2048
struct TestConstantGlobalVar {
static __constant__ int ConstantGlobalVar;
static __global__ void kernel(int* Ad) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
Ad[tid] = ConstantGlobalVar;
}
void run() {
int *A, *Ad;
A = new int[LEN];
for (unsigned i = 0; i < LEN; i++) {
A[i] = 0;
}
HIP_ASSERT(hipMalloc((void**)&Ad, SIZE));
hipLaunchKernelGGL(kernel, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0, Ad);
HIP_ASSERT(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost));
for (unsigned i = 0; i < LEN; i++) {
assert(123 == A[i]);
}
}
};
__constant__ int TestConstantGlobalVar::ConstantGlobalVar = 123;
struct TestGlobalArray {
static __device__ int GlobalArray[LEN];
static __global__ void kernelWrite() {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
GlobalArray[tid] = tid;
}
static __global__ void kernelRead(int* Ad) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
Ad[tid] = GlobalArray[tid];
}
void run() {
int *A, *Ad;
A = new int[LEN];
for (unsigned i = 0; i < LEN; i++) {
A[i] = 0;
}
HIP_ASSERT(hipMalloc((void**)&Ad, SIZE));
hipLaunchKernelGGL(kernelWrite, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0);
hipLaunchKernelGGL(kernelRead, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0, Ad);
HIP_ASSERT(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost));
for (unsigned i = 0; i < LEN; i++) {
assert(i == A[i]);
}
}
};
__device__ int TestGlobalArray::GlobalArray[LEN];
int main() {
TestConstantGlobalVar().run();
TestGlobalArray().run();
passed();
}
tests/src/runtimeApi/event/record_event.cpp
+2 -2
Ver fichero
@@ -170,8 +170,8 @@ void runTests(int64_t numElements) {
// for (int waitStart=0; waitStart<2; waitStart++) {
for (int waitStart = 1; waitStart >= 0; waitStart--) {
unsigned W = waitStart ? 0x1000 : 0;
test(W | 0x01, C_d, C_h, numElements, 0, waitStart, syncNone);
test(W | 0x02, C_d, C_h, numElements, stream, waitStart, syncNone);
test(W | 0x01, C_d, C_h, numElements, 0, 0, syncNone);
test(W | 0x02, C_d, C_h, numElements, stream, 0, syncNone);
test(W | 0x04, C_d, C_h, numElements, 0, waitStart, syncStream);
test(W | 0x08, C_d, C_h, numElements, stream, waitStart, syncStream);
test(W | 0x10, C_d, C_h, numElements, 0, waitStart, syncStopEvent);