AI/rocm-systems
2
0
Vork 0
Code Kwesties Pull-aanvragen Actions Paketten Projecten Publicaties Wiki Activiteit

Merge 'master' into 'amd-master'

Bladeren bron 
Change-Id: I143645ea2bd2c543f0d7ed736b4515c39ee2832b
This commit is contained in:
Jenkins
2019-07-18 05:22:09 -04:00
bovenliggende 174c2ca1f5 38254caf7a
commit 5575d2e52e
23 gewijzigde bestanden met toevoegingen van 171 en 218 verwijderingen
Download Patch-bestand Download Diff-bestand Expand all files Collapse all files
bin/hipify-perl
+6
Bestand weergeven
@@ -321,6 +321,12 @@ while (@ARGV) {
$ft{'err'} += s/\bcudaDevAttrMaxSharedMemoryPerMultiprocessor\b/hipDeviceAttributeMaxSharedMemoryPerMultiprocessor/g;
$ft{'err'} += s/\bcudaDevAttrMemoryClockRate\b/hipDeviceAttributeMemoryClockRate/g;
$ft{'err'} += s/\bcudaDevAttrGlobalMemoryBusWidth\b/hipDeviceAttributeMemoryBusWidth/g;
$ft{'err'} += s/\bcudaDevAttrMaxTexture1DWidth\b/hipDeviceAttributeMaxTexture1DWidth/g;
$ft{'err'} += s/\bcudaDevAttrMaxTexture2DWidth\b/hipDeviceAttributeMaxTexture2DWidth/g;
$ft{'err'} += s/\bcudaDevAttrMaxTexture2DHeight\b/hipDeviceAttributeMaxTexture2DHeight/g;
$ft{'err'} += s/\bcudaDevAttrMaxTexture3DWidth\b/hipDeviceAttributeMaxTexture3DWidth/g;
$ft{'err'} += s/\bcudaDevAttrMaxTexture3DHeight\b/hipDeviceAttributeMaxTexture3DHeight/g;
$ft{'err'} += s/\bcudaDevAttrMaxTexture3DDepth\b/hipDeviceAttributeMaxTexture3DDepth/g;
#--------
$ft{'dev'} += s/\bcudaDeviceAttr\b/hipDeviceAttribute_t/g;
$ft{'dev'} += s/\bcudaDeviceGetAttribute\b/hipDeviceGetAttribute/g;
docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md
+6 -6
Bestand weergeven
@@ -519,12 +519,12 @@
| 18 |*`cudaDevAttrIntegrated`* |*`hipDeviceAttributeIntegrated`* |
| 19 |*`cudaDevAttrCanMapHostMemory`* | |
| 20 |*`cudaDevAttrComputeMode`* |*`hipDeviceAttributeComputeMode`* |
| 21 |*`cudaDevAttrMaxTexture1DWidth`* | |
| 22 |*`cudaDevAttrMaxTexture2DWidth`* | |
| 23 |*`cudaDevAttrMaxTexture2DHeight`* | |
| 24 |*`cudaDevAttrMaxTexture3DWidth`* | |
| 25 |*`cudaDevAttrMaxTexture3DHeight`* | |
| 26 |*`cudaDevAttrMaxTexture3DDepth`* | |
| 21 |*`cudaDevAttrMaxTexture1DWidth`* |*`hipDeviceAttributeMaxTexture1DWidth`* |
| 22 |*`cudaDevAttrMaxTexture2DWidth`* |*`hipDeviceAttributeMaxTexture2DWidth`* |
| 23 |*`cudaDevAttrMaxTexture2DHeight`* |*`hipDeviceAttributeMaxTexture2DHeight`* |
| 24 |*`cudaDevAttrMaxTexture3DWidth`* |*`hipDeviceAttributeMaxTexture3DWidth`* |
| 25 |*`cudaDevAttrMaxTexture3DHeight`* |*`hipDeviceAttributeMaxTexture3DHeight`* |
| 26 |*`cudaDevAttrMaxTexture3DDepth`* |*`hipDeviceAttributeMaxTexture3DDepth`* |
| 27 |*`cudaDevAttrMaxTexture2DLayeredWidth`* | |
| 28 |*`cudaDevAttrMaxTexture2DLayeredHeight`* | |
| 29 |*`cudaDevAttrMaxTexture2DLayeredLayers`* | |
hipify-clang/src/CUDA2HIP_Driver_API_types.cpp
+6 -6
Bestand weergeven
@@ -328,17 +328,17 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
// cudaDevAttrComputeMode
{"CU_DEVICE_ATTRIBUTE_COMPUTE_MODE", {"hipDeviceAttributeComputeMode", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 20
// cudaDevAttrMaxTexture1DWidth
{"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_WIDTH", {"hipDeviceAttributeMaxTexture1DWidth", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 21
{"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_WIDTH", {"hipDeviceAttributeMaxTexture1DWidth", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 21
// cudaDevAttrMaxTexture2DWidth
{"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_WIDTH", {"hipDeviceAttributeMaxTexture2DWidth", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 22
{"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_WIDTH", {"hipDeviceAttributeMaxTexture2DWidth", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 22
// cudaDevAttrMaxTexture2DHeight
{"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_HEIGHT", {"hipDeviceAttributeMaxTexture2DHeight", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 23
{"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_HEIGHT", {"hipDeviceAttributeMaxTexture2DHeight", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 23
// cudaDevAttrMaxTexture3DWidth
{"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_WIDTH", {"hipDeviceAttributeMaxTexture3DWidth", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 24
{"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_WIDTH", {"hipDeviceAttributeMaxTexture3DWidth", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 24
// cudaDevAttrMaxTexture3DHeight
{"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_HEIGHT", {"hipDeviceAttributeMaxTexture3DHeight", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 25
{"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_HEIGHT", {"hipDeviceAttributeMaxTexture3DHeight", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 25
// cudaDevAttrMaxTexture3DDepth
{"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_DEPTH", {"hipDeviceAttributeMaxTexture3DDepth", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 26
{"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_DEPTH", {"hipDeviceAttributeMaxTexture3DDepth", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 26
// cudaDevAttrMaxTexture2DLayeredWidth
{"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_WIDTH", {"hipDeviceAttributeMaxTexture2DLayeredWidth", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 27
// cudaDevAttrMaxTexture2DLayeredHeight
hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp
+6 -6
Bestand weergeven
@@ -255,17 +255,17 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_TYPE_NAME_MAP {
// CU_DEVICE_ATTRIBUTE_COMPUTE_MODE
{"cudaDevAttrComputeMode", {"hipDeviceAttributeComputeMode", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 20
// CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_WIDTH
{"cudaDevAttrMaxTexture1DWidth", {"hipDeviceAttributeMaxTexture1DWidth", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 21
{"cudaDevAttrMaxTexture1DWidth", {"hipDeviceAttributeMaxTexture1DWidth", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 21
// CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_WIDTH
{"cudaDevAttrMaxTexture2DWidth", {"hipDeviceAttributeMaxTexture2DWidth", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 22
{"cudaDevAttrMaxTexture2DWidth", {"hipDeviceAttributeMaxTexture2DWidth", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 22
// CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_HEIGHT
{"cudaDevAttrMaxTexture2DHeight", {"hipDeviceAttributeMaxTexture2DHeight", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 23
{"cudaDevAttrMaxTexture2DHeight", {"hipDeviceAttributeMaxTexture2DHeight", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 23
// CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_WIDTH
{"cudaDevAttrMaxTexture3DWidth", {"hipDeviceAttributeMaxTexture3DWidth", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 24
{"cudaDevAttrMaxTexture3DWidth", {"hipDeviceAttributeMaxTexture3DWidth", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 24
// CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_HEIGHT
{"cudaDevAttrMaxTexture3DHeight", {"hipDeviceAttributeMaxTexture3DHeight", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 25
{"cudaDevAttrMaxTexture3DHeight", {"hipDeviceAttributeMaxTexture3DHeight", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 25
// CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_DEPTH
{"cudaDevAttrMaxTexture3DDepth", {"hipDeviceAttributeMaxTexture3DDepth", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 26
{"cudaDevAttrMaxTexture3DDepth", {"hipDeviceAttributeMaxTexture3DDepth", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 26
// CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_WIDTH
{"cudaDevAttrMaxTexture2DLayeredWidth", {"hipDeviceAttributeMaxTexture2DLayeredWidth", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 27
// CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_HEIGHT
include/hip/hcc_detail/texture_functions.h
+22 -8
Bestand weergeven
@@ -22,7 +22,6 @@ THE SOFTWARE.
#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_TEXTURE_FUNCTIONS_H
#define HIP_INCLUDE_HIP_HCC_DETAIL_TEXTURE_FUNCTIONS_H
#include <hip/hcc_detail/hip_vector_types.h>
#include <hip/hcc_detail/hip_texture_types.h>
@@ -118,8 +117,6 @@ union TData {
#define TEXTURE_RETURN_UINT return texel.u.x;
#define TEXTURE_RETURN_FLOAT return texel.f.x;
#define TEXTURE_RETURN_SIGNED return texel.i.x;
#define TEXTURE_RETURN_UNSIGNED return texel.u.x;
@@ -136,8 +133,6 @@ union TData {
#define TEXTURE_RETURN_UINT_X return make_uint1(texel.u.x);
#define TEXTURE_RETURN_FLOAT_X return make_float1(texel.f.x);
#define TEXTURE_RETURN_CHAR_XY return make_char2(texel.i.x, texel.i.y);
#define TEXTURE_RETURN_UCHAR_XY return make_uchar2(texel.u.x, texel.u.y);
@@ -150,8 +145,6 @@ union TData {
#define TEXTURE_RETURN_UINT_XY return make_uint2(texel.u.x, texel.u.y);
#define TEXTURE_RETURN_FLOAT_XY return make_float2(texel.f.x, texel.f.y);
#define TEXTURE_RETURN_CHAR_XYZW return make_char4(texel.i.x, texel.i.y, texel.i.z, texel.i.w);
#define TEXTURE_RETURN_UCHAR_XYZW return make_uchar4(texel.u.x, texel.u.y, texel.u.z, texel.u.w);
@@ -164,10 +157,30 @@ union TData {
#define TEXTURE_RETURN_UINT_XYZW return make_uint4(texel.u.x, texel.u.y, texel.u.z, texel.u.w);
#define TEXTURE_RETURN_FLOAT_XYZW return make_float4(texel.f.x, texel.f.y, texel.f.z, texel.f.w);
#define HIP_AD_FORMAT_NOT_INITIALIZED 0
#define TEXTURE_RETURN_FLOAT return (texFormatToSize[texRef.format] == 1)? texel.f.x : (float)texel.u.x/texFormatToSize[texRef.format];
#define TEXTURE_RETURN_FLOAT_X return (texFormatToSize[texRef.format] == 1)? make_float1(texel.f.x) : make_float1((float)texel.u.x/texFormatToSize[texRef.format]);
#define TEXTURE_RETURN_FLOAT_XY return (texFormatToSize[texRef.format] == 1)? make_float2(texel.f.x, texel.f.y) : make_float2((float)texel.u.x/texFormatToSize[texRef.format], (float)texel.u.y/texFormatToSize[texRef.format]);
#define TEXTURE_RETURN_FLOAT_XYZW return (texFormatToSize[texRef.format] == 1)? make_float4(texel.f.x, texel.f.y, texel.f.z, texel.f.w) : make_float4((float)texel.u.x/texFormatToSize[texRef.format], (float)texel.u.y/texFormatToSize[texRef.format], (float)texel.u.z/texFormatToSize[texRef.format], (float)texel.u.w/texFormatToSize[texRef.format]) ;
extern "C" {
__device__ __constant__ static int texFormatToSize[] = {
[HIP_AD_FORMAT_NOT_INITIALIZED] = 1 ,
[HIP_AD_FORMAT_UNSIGNED_INT8] = UCHAR_MAX ,
[HIP_AD_FORMAT_UNSIGNED_INT16]= USHRT_MAX,
[HIP_AD_FORMAT_UNSIGNED_INT32]= 1 ,
[HIP_AD_FORMAT_SIGNED_INT8] = SCHAR_MAX,
[HIP_AD_FORMAT_SIGNED_INT16] = SHRT_MAX,
[HIP_AD_FORMAT_SIGNED_INT32] = 1 ,
[HIP_AD_FORMAT_HALF] = 1 ,
[HIP_AD_FORMAT_FLOAT] = 1
};
__device__
__hip_float4_vector_value_type __ockl_image_sample_1D(
unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s,
@@ -3863,6 +3876,7 @@ __TEXTURE_FUNCTIONS_DECL__ float tex1D(texture<float, texType, mode> texRef,
TEXTURE_RETURN_FLOAT;
}
//////
template <int texType, enum hipTextureReadMode mode>
__TEXTURE_FUNCTIONS_DECL__ float tex1D(texture<float, texType, mode> texRef, float x) {
TEXTURE_REF_PARAMETERS_INIT;
include/hip/hip_runtime_api.h
+12 -1
Bestand weergeven
@@ -116,7 +116,10 @@ typedef struct hipDeviceProp_t {
int gcnArch; ///< AMD GCN Arch Value. Eg: 803, 701
int integrated; ///< APU vs dGPU
int cooperativeLaunch; ///< HIP device supports cooperative launch
int cooperativeMultiDeviceLaunch; ///< HIP device supports cooperative launch on multiple devices
int cooperativeMultiDeviceLaunch; ///< HIP device supports cooperative launch on multiple devices
int maxTexture1D; ///< Maximum number of elements in 1D images
int maxTexture2D[2]; ///< Maximum dimensions (width, height) of 2D images, in image elements
int maxTexture3D[3]; ///< Maximum dimensions (width, height, depth) of 3D images, in image elements
} hipDeviceProp_t;
@@ -295,6 +298,14 @@ 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
hipDeviceAttributeMaxTexture3DWidth, ///< Maximum dimension width of 3D images in image elements
hipDeviceAttributeMaxTexture3DHeight, ///< Maximum dimensions height of 3D images in image elements
hipDeviceAttributeMaxTexture3DDepth ///< Maximum dimensions depth of 3D images in image elements
} hipDeviceAttribute_t;
enum hipComputeMode {
include/hip/nvcc_detail/hip_runtime_api.h
+25
Bestand weergeven
@@ -772,6 +772,13 @@ inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t* p_prop, int dev
p_prop->maxSharedMemoryPerMultiProcessor = cdprop.sharedMemPerMultiprocessor;
p_prop->isMultiGpuBoard = cdprop.isMultiGpuBoard;
p_prop->maxTexture1D = cdprop.maxTexture1D;
p_prop->maxTexture2D[0] = cdprop.maxTexture2D[0];
p_prop->maxTexture2D[1] = cdprop.maxTexture2D[1];
p_prop->maxTexture3D[0] = cdprop.maxTexture3D[0];
p_prop->maxTexture3D[1] = cdprop.maxTexture3D[1];
p_prop->maxTexture3D[2] = cdprop.maxTexture3D[2];
return hipCUDAErrorTohipError(cerror);
}
@@ -858,6 +865,24 @@ inline static hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t att
case hipDeviceAttributeIntegrated:
cdattr = cudaDevAttrIntegrated;
break;
case hipDeviceAttributeMaxTexture1DWidth:
cdattr = cudaDevAttrMaxTexture1DWidth;
break;
case hipDeviceAttributeMaxTexture2DWidth:
cdattr = cudaDevAttrMaxTexture2DWidth;
break;
case hipDeviceAttributeMaxTexture2DHeight:
cdattr = cudaDevAttrMaxTexture2DHeight;
break;
case hipDeviceAttributeMaxTexture3DWidth:
cdattr = cudaDevAttrMaxTexture3DWidth;
break;
case hipDeviceAttributeMaxTexture3DHeight:
cdattr = cudaDevAttrMaxTexture3DHeight;
break;
case hipDeviceAttributeMaxTexture3DDepth:
cdattr = cudaDevAttrMaxTexture3DDepth;
break;
default:
cerror = cudaErrorInvalidValue;
break;
samples/1_Utils/hipCommander/c.cmd
+2 -2
Bestand weergeven
@@ -1,3 +1,3 @@
loop,1000; H2D; NullKernel; D2H; endloop;
loop(1000); H2D; NullKernel; D2H; endloop;
streamsync;
printTiming, 1000
printTiming(1000)
samples/1_Utils/hipCommander/classic.cmd
+1 -1
Bestand weergeven
@@ -1 +1 @@
H2D; NullKernel, D2H, streamsync
H2D; NullKernel; D2H; streamsync
samples/1_Utils/hipCommander/hipCommander.cpp
+10 -154
Bestand weergeven
@@ -7,11 +7,6 @@
#include <typeinfo>
#include <hip/hip_runtime.h>
#ifdef __HIP_PLATFORM_HCC__
#include <elf.h>
#include <hsa/hsa.h>
#include <hc.hpp>
#endif
#include <sys/time.h>
@@ -283,100 +278,6 @@ class Command {
#define FILENAME "nullkernel.hsaco"
#define KERNEL_NAME "NullKernel"
#ifdef __HIP_PLATFORM_HCC__
//=================================================================================================
// Use Aql to launch the NULL kernel.
class AqlKernelCommand : public Command {
public:
AqlKernelCommand(CommandStream* cmdStream, const std::vector<std::string> args)
: Command(cmdStream, args) {
hc::accelerator_view* av;
HIPCHECK(hipHccGetAcceleratorView(cmdStream->currentStream(), &av));
hc::accelerator acc = av->get_accelerator();
hsa_region_t systemRegion = *(hsa_region_t*)acc.get_hsa_am_system_region();
_hsaAgent = *(hsa_agent_t*)acc.get_hsa_agent();
std::ifstream file(FILENAME, std::ios::binary | std::ios::ate);
std::streamsize fsize = file.tellg();
file.seekg(0, std::ios::beg);
std::vector<char> buffer(fsize);
if (file.read(buffer.data(), fsize)) {
uint64_t elfSize = ElfSize(&buffer[0]);
assert(fsize == elfSize);
// TODO - replace module load code with explicit module load and unload.
hipModule_t module;
HIPCHECK(hipModuleLoadData(&module, &buffer[0]));
HIPCHECK(hipModuleGetFunction(&_function, module, KERNEL_NAME));
} else {
failed("could not open code object '%s'\n", FILENAME);
}
};
~AqlKernelCommand(){};
void run() override {
#define LEN 64
uint32_t len = LEN;
uint32_t one = 1;
float* Ad = NULL;
size_t argSize = 36;
char argBuffer[argSize];
*(uint32_t*)(&argBuffer[0]) = len;
*(uint32_t*)(&argBuffer[4]) = one;
*(uint32_t*)(&argBuffer[8]) = one;
*(uint32_t*)(&argBuffer[12]) = len;
*(uint32_t*)(&argBuffer[16]) = one;
*(uint32_t*)(&argBuffer[20]) = one;
*(float**)(&argBuffer[24]) = Ad; // Ad pointer argument
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0],
HIP_LAUNCH_PARAM_BUFFER_SIZE, &argSize, HIP_LAUNCH_PARAM_END};
hipModuleLaunchKernel(_function, len, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config);
};
public:
hsa_queue_t _hsaQueue;
hsa_agent_t _hsaAgent;
hipFunction_t _function;
private:
static uint64_t ElfSize(const void* emi) {
const Elf64_Ehdr* ehdr = (const Elf64_Ehdr*)emi;
const Elf64_Shdr* shdr = (const Elf64_Shdr*)((char*)emi + ehdr->e_shoff);
uint64_t max_offset = ehdr->e_shoff;
uint64_t total_size = max_offset + ehdr->e_shentsize * ehdr->e_shnum;
for (uint16_t i = 0; i < ehdr->e_shnum; ++i) {
uint64_t cur_offset = static_cast<uint64_t>(shdr[i].sh_offset);
if (max_offset < cur_offset) {
max_offset = cur_offset;
total_size = max_offset;
if (SHT_NOBITS != shdr[i].sh_type) {
total_size += static_cast<uint64_t>(shdr[i].sh_size);
}
}
}
return total_size;
}
};
#endif
//=================================================================================================
// HCC optimizes away fully NULL kernel calls, so run one that is nearly null:
class ModuleKernelCommand : public Command {
@@ -391,26 +292,18 @@ class ModuleKernelCommand : public Command {
void run() override {
#define LEN 64
uint32_t len = LEN;
uint32_t one = 1;
float *X = NULL;
HIPCHECK(hipMalloc((void**)&X, sizeof(float)));
struct {
float *Ad;
}args;
args.Ad = X;
size_t argSize = sizeof(args);
float* Ad = NULL;
size_t argSize = 36;
char argBuffer[argSize];
*(uint32_t*)(&argBuffer[0]) = len;
*(uint32_t*)(&argBuffer[4]) = one;
*(uint32_t*)(&argBuffer[8]) = one;
*(uint32_t*)(&argBuffer[12]) = len;
*(uint32_t*)(&argBuffer[16]) = one;
*(uint32_t*)(&argBuffer[20]) = one;
*(float**)(&argBuffer[24]) = Ad; // Ad pointer argument
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0],
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
HIP_LAUNCH_PARAM_BUFFER_SIZE, &argSize, HIP_LAUNCH_PARAM_END};
hipModuleLaunchKernel(_function, len, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config);
hipModuleLaunchKernel(_function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config);
};
@@ -447,33 +340,6 @@ class KernelCommand : public Command {
hipStream_t _stream;
};
#ifdef __HIP_PLATFORM_HCC__
//=================================================================================================
class PfeCommand : public Command {
public:
PfeCommand(CommandStream* cmdStream, const std::vector<std::string> args,
hipStream_t stream = 0)
: Command(cmdStream, args) {
HIPCHECK(hipHccGetAcceleratorView(stream, &_av));
}
~PfeCommand() {}
void run() override {
static const int gridX = 64;
static const int groupX = 64;
auto cf = hc::parallel_for_each(*_av, hc::extent<1>(gridX).tile(groupX),
[=](hc::index<1>& idx) __HC__ {});
}
private:
hc::accelerator_view* _av;
};
#endif
//=================================================================================================
class CopyCommand : public Command {
enum MemType { PinnedHost, UnpinnedHost, Device };
@@ -708,10 +574,8 @@ CommandStream::CommandStream(std::string commandStreamString, int iterations)
std::vector<std::string> tokens;
tokenize(commandStreamString, ';', tokens);
std::for_each(tokens.begin(), tokens.end(), [&](const std::string s) { this->parse(s); });
setStream(0);
std::for_each(tokens.begin(), tokens.end(), [&](const std::string s) { this->parse(s); });
}
@@ -831,14 +695,6 @@ void CommandStream::parse(std::string fullCmd) {
} else if (c == "vectoraddkernel") {
cmd = new KernelCommand(cmdStream, args, KernelCommand::VectorAdd);
#ifdef __HIP_PLATFORM_HCC__
} else if (c == "nullpfe") {
cmd = new PfeCommand(cmdStream, args);
} else if (c == "aqlkernel") {
cmd = new AqlKernelCommand(cmdStream, args);
#endif
} else if (c == "devicesync") {
cmd = new DeviceSyncCommand(cmdStream, args);
samples/1_Utils/hipCommander/l2.hcm
+2 -2
Bestand weergeven
@@ -1,3 +1,3 @@
setstream,1;
setstream(1);
NullKernel; streamsync;
loop,10000; H2D; NullKernel; streamsync; endloop,1;
loop(10000); H2D; NullKernel; streamsync; endloop(1);
samples/1_Utils/hipCommander/loop.hcm
+2 -2
Bestand weergeven
@@ -1,3 +1,3 @@
loop,1000; H2D; NullKernel; D2H; endloop;
loop(1000); H2D; NullKernel; D2H; endloop;
streamsync;
printTiming, 1000
printTiming(1000)
samples/1_Utils/hipCommander/loop2.hcm
+2 -2
Bestand weergeven
@@ -1,2 +1,2 @@
setstream,1;
loop,1000; NullKernel; syncstream; endloop,1,
setstream(1);
loop(1000); NullKernel; streamsync; endloop(1);
samples/1_Utils/hipCommander/nullkernel.hsaco
BIN
Bestand weergeven
Binair bestand niet weergegeven.
samples/1_Utils/hipCommander/perf/latency_hostsync.hcm
+4 -4
Bestand weergeven
@@ -1,8 +1,8 @@
setstream,1;
setstream(1);
NullKernel; streamsync;
loop,100000; NullKernel; streamsync; endloop,1;
loop(100000); NullKernel; streamsync; endloop(1);
loop,100000; H2D; streamsync; NullKernel; streamsync; endloop,1;
loop(100000); H2D; streamsync; NullKernel; streamsync; endloop(1);
loop,100000; H2D; streamsync; NullKernel; streamsync; D2H; streamsync; endloop,1;
loop(100000); H2D; streamsync; NullKernel; streamsync; D2H; streamsync; endloop(1);
samples/1_Utils/hipCommander/perf/latency_nosync.hcm
+4 -4
Bestand weergeven
@@ -1,5 +1,5 @@
setstream,1;
setstream(1);
NullKernel; streamsync;
loop,100000; NullKernel; streamsync; endloop,1;
loop,100000; H2D; NullKernel; streamsync; endloop,1;
loop,100000; H2D; NullKernel; D2H; streamsync; endloop,1;
loop(100000); NullKernel; streamsync; endloop(1);
loop(100000); H2D; NullKernel; streamsync; endloop(1);
loop(100000); H2D; NullKernel; D2H; streamsync; endloop(1);
samples/1_Utils/hipCommander/perf/latency_nullstream.hcm
+4 -4
Bestand weergeven
@@ -1,7 +1,7 @@
setstream,0;
setstream(0);
NullKernel; streamsync;
loop,100000; NullKernel; streamsync; endloop,1;
loop(100000); NullKernel; streamsync; endloop(1);
loop,100000; H2D; NullKernel; streamsync; endloop,1;
loop(100000); H2D; NullKernel; streamsync; endloop(1);
loop,100000; H2D; NullKernel; D2H; streamsync; endloop,1;
loop(100000); H2D; NullKernel; D2H; streamsync; endloop(1);
samples/1_Utils/hipCommander/perf/modulelaunch_latency.hcm
+1 -1
Bestand weergeven
@@ -1,5 +1,5 @@
setstream(1);
NullKernel; streamsync;
loop(100); ModuleKernel; streamsync; endloop(1);
loop(100); AqlKernel; streamsync; endloop(1);
loop(100); streamsync; endloop(1);
loop(3000); NullKernel; streamsync; endloop(1);
samples/1_Utils/hipCommander/setstream.hcm
+2 -2
Bestand weergeven
@@ -1,3 +1,3 @@
setstream,1;
setstream,2; H2D; NullKernel; D2H;
setstream(1);
setstream(2); H2D; NullKernel; D2H;
streamsync
src/hip_device.cpp
+18
Bestand weergeven
@@ -276,6 +276,24 @@ hipError_t ihipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device
case hipDeviceAttributeIntegrated:
*pi = prop->integrated;
break;
case hipDeviceAttributeMaxTexture1DWidth:
*pi = prop->maxTexture1D;
break;
case hipDeviceAttributeMaxTexture2DWidth:
*pi = prop->maxTexture2D[0];
break;
case hipDeviceAttributeMaxTexture2DHeight:
*pi = prop->maxTexture2D[1];
break;
case hipDeviceAttributeMaxTexture3DWidth:
*pi = prop->maxTexture3D[0];
break;
case hipDeviceAttributeMaxTexture3DHeight:
*pi = prop->maxTexture3D[1];
break;
case hipDeviceAttributeMaxTexture3DDepth:
*pi = prop->maxTexture3D[2];
break;
default:
e = hipErrorInvalidValue;
break;
src/hip_hcc.cpp
+28 -12
Bestand weergeven
@@ -43,7 +43,7 @@ THE SOFTWARE.
#include <hc.hpp>
#include <hc_am.hpp>
#include "hsa/hsa_ext_amd.h"
#include "hsa/hsa_ext_image.h"
#include "hip/hip_runtime.h"
#include "hip_hcc_internal.h"
#include "trace_helper.h"
@@ -913,6 +913,18 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop) {
if(agent_profile == HSA_PROFILE_FULL) {
prop->integrated = 1;
}
err = hsa_agent_get_info(_hsaAgent, (hsa_agent_info_t)HSA_EXT_AGENT_INFO_IMAGE_1D_MAX_ELEMENTS,
&prop->maxTexture1D);
DeviceErrorCheck(err);
err = hsa_agent_get_info(_hsaAgent, (hsa_agent_info_t)HSA_EXT_AGENT_INFO_IMAGE_2D_MAX_ELEMENTS,
prop->maxTexture2D);
DeviceErrorCheck(err);
err = hsa_agent_get_info(_hsaAgent, (hsa_agent_info_t)HSA_EXT_AGENT_INFO_IMAGE_3D_MAX_ELEMENTS,
prop->maxTexture3D);
DeviceErrorCheck(err);
return e;
}
@@ -1399,7 +1411,7 @@ void ihipInit() {
// Make sure the hip visible devices are within the deviceCnt range
for (int i = 0; i < g_hip_visible_devices.size(); i++) {
if (g_hip_visible_devices[i] >= deviceCnt) {
if ((g_hip_visible_devices[i] >= deviceCnt) ||(g_hip_visible_devices[i] < 0)){
// Make sure any DeviceID after invalid DeviceID will be erased.
g_hip_visible_devices.resize(i);
break;
@@ -1415,17 +1427,21 @@ void ihipInit() {
g_deviceArray = new ihipDevice_t*[deviceCnt];
g_deviceCnt = 0;
for (int i = 0; i < accs.size(); i++) {
// check if the device id is included in the HIP_VISIBLE_DEVICES env variable
if (!accs[i].get_is_emulated()) {
if (std::find(g_hip_visible_devices.begin(), g_hip_visible_devices.end(), (i - 1)) ==
g_hip_visible_devices.end() &&
g_visible_device) {
// If device is not in visible devices list, ignore
continue;
if(g_visible_device) {
for (int i = 0; i < g_hip_visible_devices.size(); i++) {
int devIndex = g_hip_visible_devices[i];
if (!accs[devIndex+1].get_is_emulated()) {
g_deviceArray[g_deviceCnt] = new ihipDevice_t(g_deviceCnt, deviceCnt, accs[devIndex+1]);
g_deviceCnt++;
}
}
}else {
for (int i = 0; i < accs.size(); i++) {
if (!accs[i].get_is_emulated()) {
g_deviceArray[g_deviceCnt] = new ihipDevice_t(g_deviceCnt, deviceCnt, accs[i]);
g_deviceCnt++;
}
g_deviceArray[g_deviceCnt] = new ihipDevice_t(g_deviceCnt, deviceCnt, accs[i]);
g_deviceCnt++;
}
}
src/hip_module.cpp
+1 -1
Bestand weergeven
@@ -993,7 +993,7 @@ hipError_t ihipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* bloc
tmp_min = min(maxWavesSGPRSLimited, tmp_min);
activeWavefronts = min(maxWavesVGPRSLimited, tmp_min);
if (maxActivWaves < activeWavefronts) {
if (maxActivWaves <= activeWavefronts) {
maxActivWaves = activeWavefronts;
maxWavefronts = wavefrontsPerWG;
}
tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp
+7
Bestand weergeven
@@ -109,5 +109,12 @@ int main(int argc, char* argv[]) {
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxSharedMemoryPerMultiprocessor,
props.maxSharedMemoryPerMultiProcessor));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeIntegrated, props.integrated));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture1DWidth, props.maxTexture1D));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture2DWidth, props.maxTexture2D[0]));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture2DHeight, props.maxTexture2D[1]));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture3DWidth, props.maxTexture3D[0]));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture3DHeight, props.maxTexture3D[1]));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture3DDepth, props.maxTexture3D[2]));
passed();
};