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

Merge 'master' into 'amd-master'

Browse Source 
Change-Id: I683a764c0546fe579f5ad32863c6c94843d3e7e8


[ROCm/clr commit: 051c9a47ac]
This commit is contained in:
Jenkins
2019-08-10 20:58:39 -04:00
parent 50b3ddee6c f57dbcc4a0
commit c0c20b91cd
17 changed files with 1327 additions and 95 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/CMakeLists.txt
+11 -4
View File
@@ -28,16 +28,23 @@ list(GET VERSION_LIST 0 HIP_VERSION_MAJOR)
list(GET VERSION_LIST 1 HIP_VERSION_MINOR) list(GET VERSION_LIST 1 HIP_VERSION_MINOR)
# get date information based on UTC # get date information based on UTC
# use the last two digits of year + week number + day in the week as HIP_VERSION_PATCH # use the last two digits of year + week number + day in the week as HIP_VERSION_GITDATE
# use the commit date, instead of build date # use the commit date, instead of build date
# add xargs to remove strange trailing newline character # add xargs to remove strange trailing newline character
execute_process(COMMAND git show -s --format=@%ct execute_process(COMMAND git show -s --format=@%ct
COMMAND xargs COMMAND xargs
COMMAND date -f - --utc +%y%U%w COMMAND date -f - --utc +%y%U%w
WORKING_DIRECTORY ${PROJECT_SOURCE_DIR} WORKING_DIRECTORY ${PROJECT_SOURCE_DIR}
OUTPUT_VARIABLE HIP_VERSION_PATCH OUTPUT_VARIABLE HIP_VERSION_GITDATE
OUTPUT_STRIP_TRAILING_WHITESPACE) OUTPUT_STRIP_TRAILING_WHITESPACE)
# get commit short hash
execute_process(COMMAND git rev-parse --short HEAD
WORKING_DIRECTORY ${PROJECT_SOURCE_DIR}
OUTPUT_VARIABLE HIP_VERSION_GITHASH
OUTPUT_STRIP_TRAILING_WHITESPACE)
set(HIP_VERSION_PATCH ${HIP_VERSION_GITDATE}-${HIP_VERSION_GITHASH})
set(HIP_VERSION ${HIP_VERSION_MAJOR}.${HIP_VERSION_MINOR}.${HIP_VERSION_PATCH}) set(HIP_VERSION ${HIP_VERSION_MAJOR}.${HIP_VERSION_MINOR}.${HIP_VERSION_PATCH})
add_to_config(_versionInfo HIP_VERSION_MAJOR) add_to_config(_versionInfo HIP_VERSION_MAJOR)
add_to_config(_versionInfo HIP_VERSION_MINOR) add_to_config(_versionInfo HIP_VERSION_MINOR)
@@ -218,7 +225,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
endif() endif()
# Add HIP_VERSION to CMAKE_<LANG>_FLAGS # Add HIP_VERSION to CMAKE_<LANG>_FLAGS
set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -DHIP_VERSION_MAJOR=${HIP_VERSION_MAJOR} -DHIP_VERSION_MINOR=${HIP_VERSION_MINOR} -DHIP_VERSION_PATCH=${HIP_VERSION_PATCH}") set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -DHIP_VERSION_MAJOR=${HIP_VERSION_MAJOR} -DHIP_VERSION_MINOR=${HIP_VERSION_MINOR} -DHIP_VERSION_PATCH=${HIP_VERSION_GITDATE}")
# Add remaining flags # Add remaining flags
set(HCC_CXX_FLAGS "-hc -fno-gpu-rdc --amdgpu-target=gfx803 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 ") set(HCC_CXX_FLAGS "-hc -fno-gpu-rdc --amdgpu-target=gfx803 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 ")
@@ -370,7 +377,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
write_basic_package_version_file( write_basic_package_version_file(
${CMAKE_CURRENT_BINARY_DIR}/hip-config-version.cmake ${CMAKE_CURRENT_BINARY_DIR}/hip-config-version.cmake
VERSION "${HIP_VERSION}" VERSION "${HIP_VERSION_MAJOR}.${HIP_VERSION_MINOR}.${HIP_VERSION_GITDATE}"
COMPATIBILITY SameMajorVersion COMPATIBILITY SameMajorVersion
) )
install( install(
projects/clr/hipamd/bin/hipcc
+2 -1
View File
@@ -102,6 +102,7 @@ sub delete_temp_dirs {
$HIP_PLATFORM= `$HIP_PATH/bin/hipconfig --platform` // "hcc"; $HIP_PLATFORM= `$HIP_PATH/bin/hipconfig --platform` // "hcc";
$HIP_VERSION= `$HIP_PATH/bin/hipconfig --version`; $HIP_VERSION= `$HIP_PATH/bin/hipconfig --version`;
($HIP_VERSION_MAJOR, $HIP_VERSION_MINOR, $HIP_VERSION_PATCH) = split(/\./, $HIP_VERSION); ($HIP_VERSION_MAJOR, $HIP_VERSION_MINOR, $HIP_VERSION_PATCH) = split(/\./, $HIP_VERSION);
my ($HIP_VERSION_GITDATE, $HIP_VERSION_GITHASH) = split(/-/, $HIP_VERSION_PATCH);
$HIP_COMPILER= $hipConfig{'HIP_COMPILER'}; $HIP_COMPILER= $hipConfig{'HIP_COMPILER'};
$HIP_RUNTIME= $hipConfig{'HIP_RUNTIME'}; $HIP_RUNTIME= $hipConfig{'HIP_RUNTIME'};
@@ -323,7 +324,7 @@ if ($HIP_PLATFORM eq "clang") {
} }
# Add paths to common HIP includes: # Add paths to common HIP includes:
$HIPCXXFLAGS .= " -isystem $HIP_INCLUDE_PATH -DHIP_VERSION_MAJOR=$HIP_VERSION_MAJOR -DHIP_VERSION_MINOR=$HIP_VERSION_MINOR -DHIP_VERSION_PATCH=$HIP_VERSION_PATCH" ; $HIPCXXFLAGS .= " -isystem $HIP_INCLUDE_PATH -DHIP_VERSION_MAJOR=$HIP_VERSION_MAJOR -DHIP_VERSION_MINOR=$HIP_VERSION_MINOR -DHIP_VERSION_PATCH=$HIP_VERSION_GITDATE" ;
my $compileOnly = 0; my $compileOnly = 0;
my $needCXXFLAGS = 0; # need to add CXX flags to compile step my $needCXXFLAGS = 0; # need to add CXX flags to compile step
projects/clr/hipamd/bin/hipconfig
+2 -2
View File
@@ -1,7 +1,7 @@
#!/usr/bin/perl -w #!/usr/bin/perl -w
$HIP_BASE_VERSION_MAJOR = "1"; $HIP_BASE_VERSION_MAJOR = "2";
$HIP_BASE_VERSION_MINOR = "5"; $HIP_BASE_VERSION_MINOR = "8";
# Need perl > 5.10 to use logic-defined or # Need perl > 5.10 to use logic-defined or
use 5.006; use v5.10.1; use 5.006; use v5.10.1;
projects/clr/hipamd/docs/markdown/CURAND_API_supported_by_HIP.md
+18 -10
View File
@@ -93,16 +93,24 @@
| 11 |*`CURAND_3RD`* | | | 11 |*`CURAND_3RD`* | |
| 12 |*`CURAND_DEFINITION`* | | | 12 |*`CURAND_DEFINITION`* | |
| 13 |*`CURAND_POISSON`* | | | 13 |*`CURAND_POISSON`* | |
| struct | `curandStateMtgp32_t` | `hiprandStateMtgp32_t` | | struct | `curandStateMtgp32` | `hiprandStateMtgp32` |
| struct | `curandStateScrambledSobol64_t` | | | typedef | `curandStateMtgp32_t` | `hiprandStateMtgp32_t` |
| struct | `curandStateSobol64_t` | | | struct | `curandStateScrambledSobol64` | |
| struct | `curandStateScrambledSobol32_t` | | | typedef | `curandStateScrambledSobol64_t` | |
| struct | `curandStateSobol32_t` | `hiprandStateSobol32_t` | | struct | `curandStateSobol64` | |
| struct | `curandStateMRG32k3a_t` | `hiprandStateMRG32k3a_t` | | typedef | `curandStateSobol64_t` | |
| struct | `curandStatePhilox4_32_10_t` | `hiprandStatePhilox4_32_10_t` | | struct | `curandStateScrambledSobol32` | |
| struct | `curandStateXORWOW_t` | `hiprandStateXORWOW_t` | | typedef | `curandStateScrambledSobol32_t` | |
| struct | `curandState_t` | `hiprandState_t` | | struct | `curandStateSobol32` | `hiprandStateSobol32` |
| struct | `curandState` | `hiprandState_t` | | typedef | `curandStateSobol32_t` | `hiprandStateSobol32_t` |
| struct | `curandStateMRG32k3a` | `hiprandStateMRG32k3a` |
| typedef | `curandStateMRG32k3a_t` | `hiprandStateMRG32k3a_t` |
| struct | `curandStatePhilox4_32_10` | `hiprandStatePhilox4_32_10` |
| typedef | `curandStatePhilox4_32_10_t` | `hiprandStatePhilox4_32_10_t` |
| struct | `curandStateXORWOW` | `hiprandStateXORWOW` |
| typedef | `curandStateXORWOW_t` | `hiprandStateXORWOW_t` |
| struct | `curandState` | `hiprandState` |
| typedef | `curandState_t` | `hiprandState_t` |
## **2. Host API Functions** ## **2. Host API Functions**
projects/clr/hipamd/docs/markdown/hip_porting_guide.md
+2 -2
View File
@@ -186,8 +186,8 @@ Unlike `__CUDA_ARCH__`, the `__HIP_DEVICE_COMPILE__` value is 1 or undefined, an
|Define | hcc | hip-clang | nvcc | Other (GCC, ICC, Clang, etc.) |Define | hcc | hip-clang | nvcc | Other (GCC, ICC, Clang, etc.)
|--- | --- | --- | --- |---| |--- | --- | --- | --- |---|
|HIP-related defines:| |HIP-related defines:|
|`__HIP_PLATFORM_HCC___`| Defined | Defined | Undefined | Defined if targeting hcc platform; undefined otherwise | |`__HIP_PLATFORM_HCC__`| Defined | Defined | Undefined | Defined if targeting hcc platform; undefined otherwise |
|`__HIP_PLATFORM_NVCC___`| Undefined | Undefined | Defined | Defined if targeting nvcc platform; undefined otherwise | |`__HIP_PLATFORM_NVCC__`| Undefined | Undefined | Defined | Defined if targeting nvcc platform; undefined otherwise |
|`__HIP_DEVICE_COMPILE__` | 1 if compiling for device; undefined if compiling for host | 1 if compiling for device; undefined if compiling for host |1 if compiling for device; undefined if compiling for host | Undefined |`__HIP_DEVICE_COMPILE__` | 1 if compiling for device; undefined if compiling for host | 1 if compiling for device; undefined if compiling for host |1 if compiling for device; undefined if compiling for host | Undefined
|`__HIPCC__` | Defined | Defined | Defined | Undefined |`__HIPCC__` | Defined | Defined | Defined | Undefined
|`__HIP_ARCH_*` | 0 or 1 depending on feature support (see below) |0 or 1 depending on feature support (see below) | 0 or 1 depending on feature support (see below) | 0 |`__HIP_ARCH_*` | 0 or 1 depending on feature support (see below) |0 or 1 depending on feature support (see below) | 0 or 1 depending on feature support (see below) | 0
projects/clr/hipamd/hipify-clang/src/CUDA2HIP_RAND_API_types.cpp
+9 -2
View File
@@ -53,17 +53,24 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RAND_TYPE_NAME_MAP{
{"curandDirectionVectors32_t", {"hiprandDirectionVectors32_t", "", CONV_TYPE, API_RAND}}, {"curandDirectionVectors32_t", {"hiprandDirectionVectors32_t", "", CONV_TYPE, API_RAND}},
{"curandDirectionVectors64_t", {"hiprandDirectionVectors64_t", "", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, {"curandDirectionVectors64_t", {"hiprandDirectionVectors64_t", "", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}},
// cuRAND types for Device functions // cuRAND types for Device functions
{"curandStateMtgp32", {"hiprandStateMtgp32", "", CONV_TYPE, API_RAND}},
{"curandStateMtgp32_t", {"hiprandStateMtgp32_t", "", CONV_TYPE, API_RAND}}, {"curandStateMtgp32_t", {"hiprandStateMtgp32_t", "", CONV_TYPE, API_RAND}},
{"curandStateScrambledSobol64", {"hiprandStateScrambledSobol64", "", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}},
{"curandStateScrambledSobol64_t", {"hiprandStateScrambledSobol64_t", "", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, {"curandStateScrambledSobol64_t", {"hiprandStateScrambledSobol64_t", "", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}},
{"curandStateSobol64", {"hiprandStateSobol64", "", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}},
{"curandStateSobol64_t", {"hiprandStateSobol64_t", "", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, {"curandStateSobol64_t", {"hiprandStateSobol64_t", "", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}},
{"curandStateScrambledSobol32", {"hiprandStateScrambledSobol32", "", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}},
{"curandStateScrambledSobol32_t", {"hiprandStateScrambledSobol32_t", "", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, {"curandStateScrambledSobol32_t", {"hiprandStateScrambledSobol32_t", "", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}},
{"curandStateSobol32", {"hiprandStateSobol32", "", CONV_TYPE, API_RAND}},
{"curandStateSobol32_t", {"hiprandStateSobol32_t", "", CONV_TYPE, API_RAND}}, {"curandStateSobol32_t", {"hiprandStateSobol32_t", "", CONV_TYPE, API_RAND}},
{"curandStateMRG32k3a", {"hiprandStateMRG32k3a", "", CONV_TYPE, API_RAND}},
{"curandStateMRG32k3a_t", {"hiprandStateMRG32k3a_t", "", CONV_TYPE, API_RAND}}, {"curandStateMRG32k3a_t", {"hiprandStateMRG32k3a_t", "", CONV_TYPE, API_RAND}},
{"curandStatePhilox4_32_10", {"hiprandStatePhilox4_32_10", "", CONV_TYPE, API_RAND}},
{"curandStatePhilox4_32_10_t", {"hiprandStatePhilox4_32_10_t", "", CONV_TYPE, API_RAND}}, {"curandStatePhilox4_32_10_t", {"hiprandStatePhilox4_32_10_t", "", CONV_TYPE, API_RAND}},
{"curandStateXORWOW", {"hiprandStateXORWOW", "", CONV_TYPE, API_RAND}},
{"curandStateXORWOW_t", {"hiprandStateXORWOW_t", "", CONV_TYPE, API_RAND}}, {"curandStateXORWOW_t", {"hiprandStateXORWOW_t", "", CONV_TYPE, API_RAND}},
{"curandState", {"hiprandState", "", CONV_TYPE, API_RAND}},
{"curandState_t", {"hiprandState_t", "", CONV_TYPE, API_RAND}}, {"curandState_t", {"hiprandState_t", "", CONV_TYPE, API_RAND}},
{"curandState", {"hiprandState_t", "", CONV_TYPE, API_RAND}},
// RAND function call status types (enum curandStatus) // RAND function call status types (enum curandStatus)
{"CURAND_STATUS_SUCCESS", {"HIPRAND_STATUS_SUCCESS", "", CONV_NUMERIC_LITERAL, API_RAND}}, {"CURAND_STATUS_SUCCESS", {"HIPRAND_STATUS_SUCCESS", "", CONV_NUMERIC_LITERAL, API_RAND}},
projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h
+15 -3
View File
@@ -1882,15 +1882,27 @@ hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch,
size_t height, hipMemcpyKind kind); size_t height, hipMemcpyKind kind);
/** /**
* @brief Copies memory for 2D arrays. * @brief Copies memory for 2D arrays.
* @param[in] pCopy Parameters for the memory copy * @param[in] pCopy Parameters for the memory copy
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue,
* #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection
*
* @see hipMemcpy, hipMemcpy2D, hipMemcpyToArray, hipMemcpy2DToArray, hipMemcpyFromArray,
* hipMemcpyToSymbol, hipMemcpyAsync
*/
hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy);
/**
* @brief Copies memory for 2D arrays.
* @param[in] pCopy Parameters for the memory copy
* @param[in] stream Stream to use
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue,
* #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection * #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection
* *
* @see hipMemcpy, hipMemcpy2D, hipMemcpyToArray, hipMemcpy2DToArray, hipMemcpyFromArray, * @see hipMemcpy, hipMemcpy2D, hipMemcpyToArray, hipMemcpy2DToArray, hipMemcpyFromArray,
* hipMemcpyToSymbol, hipMemcpyAsync * hipMemcpyToSymbol, hipMemcpyAsync
*/ */
hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy); hipError_t hipMemcpyParam2DAsync(const hip_Memcpy2D* pCopy, hipStream_t stream __dparm(0));
/** /**
* @brief Copies data between host and device. * @brief Copies data between host and device.
projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h
+4
View File
@@ -597,6 +597,10 @@ inline static hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) {
return hipCUResultTohipError(cuMemcpy2D(pCopy)); return hipCUResultTohipError(cuMemcpy2D(pCopy));
} }
inline static hipError_t hipMemcpyParam2DAsync(const hip_Memcpy2D* pCopy, hipStream_t stream __dparm(0)) {
return hipCUResultTohipError(cuMemcpy2DAsync(pCopy, stream));
}
inline static hipError_t hipMemcpy3D(const struct hipMemcpy3DParms *p) inline static hipError_t hipMemcpy3D(const struct hipMemcpy3DParms *p)
{ {
return hipCUDAErrorTohipError(cudaMemcpy3D(p)); return hipCUDAErrorTohipError(cudaMemcpy3D(p));
projects/clr/hipamd/src/hip_memory.cpp
+25 -5
View File
@@ -1694,10 +1694,9 @@ hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch,
return ihipLogStatus(e); return ihipLogStatus(e);
} }
hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, hipError_t ihipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
size_t height, hipMemcpyKind kind, hipStream_t stream) { size_t height, hipMemcpyKind kind, hipStream_t stream) {
HIP_INIT_SPECIAL_API(hipMemcpy2DAsync, (TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind, stream); if (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return hipErrorInvalidValue;
if (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return ihipLogStatus(hipErrorInvalidValue);
hipError_t e = hipSuccess; hipError_t e = hipSuccess;
int isLockedOrD2D = 0; int isLockedOrD2D = 0;
void *pinnedPtr=NULL; void *pinnedPtr=NULL;
@@ -1736,6 +1735,14 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
} }
} }
return e;
}
hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
size_t height, hipMemcpyKind kind, hipStream_t stream) {
HIP_INIT_SPECIAL_API(hipMemcpy2DAsync, (TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind, stream);
hipError_t e = hipSuccess;
e = ihipMemcpy2DAsync(dst, dpitch, src, spitch, width, height, kind, stream);
return ihipLogStatus(e); return ihipLogStatus(e);
} }
@@ -1744,9 +1751,22 @@ hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) {
hipError_t e = hipSuccess; hipError_t e = hipSuccess;
if (pCopy == nullptr) { if (pCopy == nullptr) {
e = hipErrorInvalidValue; e = hipErrorInvalidValue;
} } else {
e = ihipMemcpy2D(pCopy->dstArray->data, pCopy->WidthInBytes, pCopy->srcHost, pCopy->srcPitch, e = ihipMemcpy2D(pCopy->dstArray->data, pCopy->WidthInBytes, pCopy->srcHost, pCopy->srcPitch,
pCopy->WidthInBytes, pCopy->Height, hipMemcpyDefault); pCopy->WidthInBytes, pCopy->Height, hipMemcpyDefault);
}
return ihipLogStatus(e);
}
hipError_t hipMemcpyParam2DAsync(const hip_Memcpy2D* pCopy, hipStream_t stream) {
HIP_INIT_SPECIAL_API(hipMemcpyParam2DAsync, (TRACE_MCMD), pCopy, stream);
hipError_t e = hipSuccess;
if (pCopy == nullptr) {
e = hipErrorInvalidValue;
} else {
e = ihipMemcpy2DAsync(pCopy->dstArray->data, pCopy->WidthInBytes, pCopy->srcHost, pCopy->srcPitch,
pCopy->WidthInBytes, pCopy->Height, hipMemcpyDefault, stream);
}
return ihipLogStatus(e); return ihipLogStatus(e);
} }
projects/clr/hipamd/tests/hipify-clang/unit_tests/libraries/cuRAND/poisson_api_example.cu
+9 -10
View File
@@ -52,7 +52,7 @@ API_TYPE api;
int report_break; int report_break;
int cashiers_load_h[HOURS]; int cashiers_load_h[HOURS];
__constant__ int cashiers_load[HOURS]; __constant__ int cashiers_load[HOURS];
// CHECK: __global__ void setup_kernel(hiprandState_t *state) // CHECK: __global__ void setup_kernel(hiprandState *state)
__global__ void setup_kernel(curandState *state) __global__ void setup_kernel(curandState *state)
{ {
int id = threadIdx.x + blockIdx.x * blockDim.x; int id = threadIdx.x + blockIdx.x * blockDim.x;
@@ -79,15 +79,15 @@ void update_queue(int id, int min, unsigned int new_customers,
= queue_length; = queue_length;
} }
// CHECK: __global__ void simple_device_API_kernel(hiprandState_t *state, // CHECK: __global__ void simple_device_API_kernel(hiprandState *state,
__global__ void simple_device_API_kernel(curandState *state, __global__ void simple_device_API_kernel(curandState *state,
unsigned int *queue_lengths, size_t pitch) unsigned int *queue_lengths, size_t pitch)
{ {
int id = threadIdx.x + blockIdx.x * blockDim.x; int id = threadIdx.x + blockIdx.x * blockDim.x;
unsigned int new_customers; unsigned int new_customers;
unsigned int queue_length = 0; unsigned int queue_length = 0;
/* Copy state to local memory for efficiency */ /* Copy state to local memory for efficiency */
// CHECK: hiprandState_t localState = state[id]; // CHECK: hiprandState localState = state[id];
curandState localState = state[id]; curandState localState = state[id];
/* Simulate queue in time */ /* Simulate queue in time */
for(int min = 1; min <= 60 * HOURS; min++) { for(int min = 1; min <= 60 * HOURS; min++) {
@@ -120,7 +120,7 @@ __global__ void host_API_kernel(unsigned int *poisson_numbers,
queue_lengths, pitch); queue_lengths, pitch);
} }
} }
// CHECK: __global__ void robust_device_API_kernel(hiprandState_t *state, // CHECK: __global__ void robust_device_API_kernel(hiprandState *state,
// CHECK: hiprandDiscreteDistribution_t poisson_1, // CHECK: hiprandDiscreteDistribution_t poisson_1,
// CHECK: hiprandDiscreteDistribution_t poisson_2, // CHECK: hiprandDiscreteDistribution_t poisson_2,
// CHECK: hiprandDiscreteDistribution_t poisson_3, // CHECK: hiprandDiscreteDistribution_t poisson_3,
@@ -134,7 +134,7 @@ __global__ void robust_device_API_kernel(curandState *state,
unsigned int new_customers; unsigned int new_customers;
unsigned int queue_length = 0; unsigned int queue_length = 0;
/* Copy state to local memory for efficiency */ /* Copy state to local memory for efficiency */
// CHECK: hiprandState_t localState = state[id]; // CHECK: hiprandState localState = state[id];
curandState localState = state[id]; curandState localState = state[id];
/* Simulate queue in time */ /* Simulate queue in time */
/* first 3 hours */ /* first 3 hours */
@@ -165,7 +165,7 @@ __global__ void robust_device_API_kernel(curandState *state,
curand_discrete(&localState, poisson_3); curand_discrete(&localState, poisson_3);
/* Update queue */ /* Update queue */
update_queue(id, min, new_customers, queue_length, update_queue(id, min, new_customers, queue_length,
queue_lengths, pitch); queue_lengths, pitch);
} }
/* Copy state back to global memory */ /* Copy state back to global memory */
state[id] = localState; state[id] = localState;
@@ -298,12 +298,11 @@ void print_statistics(unsigned int *hostResults, size_t pitch)
} }
} }
int main(int argc, char *argv[]) int main(int argc, char *argv[])
{ {
int n; int n;
size_t pitch; size_t pitch;
// CHECK: hiprandState_t *devStates; // CHECK: hiprandState *devStates;
curandState *devStates; curandState *devStates;
unsigned int *devResults, *hostResults; unsigned int *devResults, *hostResults;
unsigned int *poisson_numbers_d; unsigned int *poisson_numbers_d;
@@ -328,7 +327,7 @@ int main(int argc, char *argv[])
/* Allocate space for prng states on device */ /* Allocate space for prng states on device */
// CHECK: CUDA_CALL(hipMalloc((void **)&devStates, 64 * 64 * // CHECK: CUDA_CALL(hipMalloc((void **)&devStates, 64 * 64 *
// CHECK: sizeof(hiprandState_t))); // CHECK: sizeof(hiprandState)));
CUDA_CALL(cudaMalloc((void **)&devStates, 64 * 64 * CUDA_CALL(cudaMalloc((void **)&devStates, 64 * 64 *
sizeof(curandState))); sizeof(curandState)));
projects/clr/hipamd/tests/src/clara/clara.hpp
+1165
View File
File diff suppressed because it is too large Load Diff
projects/clr/hipamd/tests/src/deviceLib/hipAsynchronousStreams.cpp
+2 -2
View File
@@ -108,12 +108,12 @@ typedef struct {
extern "C" __device__ __ockl_as_status_t extern "C" __device__ __ockl_as_status_t
__ockl_as_write_block(__ockl_as_stream_t ATTR_GLOBAL *stream, uchar service_id, __ockl_as_write_block(__ockl_as_stream_t ATTR_GLOBAL *stream, uchar service_id,
ulong *connection_id, const uchar *str, uint32_t len, uint64_t *connection_id, const uchar *str, uint32_t len,
uchar flags); uchar flags);
__device__ __ockl_as_status_t __device__ __ockl_as_status_t
__hip_as_write_block(__ockl_as_stream_t *stream, uchar service_id, __hip_as_write_block(__ockl_as_stream_t *stream, uchar service_id,
ulong *connection_id, const uchar *str, uint32_t len, uint64_t *connection_id, const uchar *str, uint32_t len,
uchar flags) uchar flags)
{ {
__ockl_as_stream_t ATTR_GLOBAL *gstream = __ockl_as_stream_t ATTR_GLOBAL *gstream =
projects/clr/hipamd/tests/src/hipEnvVar.cpp
+41 -45
View File
@@ -27,57 +27,56 @@ THE SOFTWARE.
#include <stdio.h> #include <stdio.h>
#include <stdlib.h> #include <stdlib.h>
#include <getopt.h>
#include <iostream> #include <iostream>
#include "clara/clara.hpp"
#include <string> #include <string>
#include "hip/hip_runtime.h" #include "hip/hip_runtime.h"
#ifdef _WIN64
#include <tchar.h>
#endif
#ifdef _WIN64
#define setenv(x,y,z) _putenv_s(x,y)
#endif
using namespace std; using namespace std;
using namespace clara;
inline clara::Parser cmdline_parser(bool& help, std::string& env, int &device, bool &retDevCnt) {
return clara::Opt{retDevCnt}
["-c"]
("total number of GPUs available") |
void usage() { clara::Help{help} |
printf(
"hipEnvVar [otpions]\n\ clara::Opt{device,"device"}
-c,\t\ttotal number of available GPUs and their pciBusID\n\ ["-d"]["--device"]
-d,\t\tselect one GPU and return its pciBusID\n\ ("select one GPU and return its pciBusID") |
-v,\t\tsend the list to HIP_VISIBLE_DEVICES env var\n\
-h,\t\tshow this help message\n\ clara::Opt{env,"Set Env Value"}
"); ["-v"]["--EnvValue"]
("send the list to HIP_VISIBLE_DEVICES env var, syntax -v=<value>");
} }
int main(int argc, char** argv) {
// string str = getenv("HIP_VISIBLE_DEVICES");
// std::cout << "The current env HIP_VISIBLE_DEVICES is"<<str << std::endl;
extern char* optarg;
extern int optind;
int c = 0;
int retDevCnt = 0, retDevInfo = 0, setEnvVar = 0;
int device = 0;
string env;
while ((c = getopt(argc, argv, "cd:v:h")) != -1) switch (c) {
case 'c':
retDevCnt = true;
break;
case 'd':
retDevInfo = true;
device = atoi(optarg);
break;
case 'v':
setEnvVar = true;
env = optarg;
break;
case 'h':
usage();
return 0;
default:
// usage();
return -1;
}
if (setEnvVar) { int main(int argc, char** argv) {
// env = "export HIP_VISIBLE_DEVICES=" + env; bool help = false;
// cout<<"The received env var is: "<<env<<endl; bool retDevCnt = false;
int c = 0;
int device = INT_MAX;
string env;
auto cmd = cmdline_parser(help, env, device, retDevCnt);
const auto r = cmd.parse(Args{argc, argv});
if (!r) { std::cout<<"Valid device must be >= 0"<<std::endl; return -1;}
if (help)
cout << cmd << endl;
if (!env.empty()) {
setenv("HIP_VISIBLE_DEVICES", env.c_str(), 1); setenv("HIP_VISIBLE_DEVICES", env.c_str(), 1);
setenv("CUDA_VISIBLE_DEVICES", env.c_str(), 1); setenv("CUDA_VISIBLE_DEVICES", env.c_str(), 1);
cout << "set env HIP_VISIBLE_DEVICES = " << env.c_str() << endl; cout << "set env HIP_VISIBLE_DEVICES = " << env.c_str() << endl;
// verify if the environment variable is set // verify if the environment variable is set
char* pPath; char* pPath;
pPath = getenv("HIP_VISIBLE_DEVICES"); pPath = getenv("HIP_VISIBLE_DEVICES");
@@ -91,23 +90,20 @@ int main(int argc, char** argv) {
int devCount = 0; int devCount = 0;
hipGetDeviceCount(&devCount); hipGetDeviceCount(&devCount);
// printf("\nTotal number of GPU devices in the system is %d\n",devCount);
if (devCount == 0) { if (devCount == 0) {
printf("No HIP enabled device\n"); printf("No HIP enabled device\n");
return -1; return -1;
} }
if (device < 0 || device > devCount - 1) { if (device != INT_MAX && (device < 0 || device > devCount - 1)) {
printf("Selected device %d is out of bound. Devices on your system are in range %d - %d\n", printf("Selected device %d is out of bound. Devices on your system are in range %d - %d\n",
device, 0, devCount - 1); device, 0, devCount - 1);
return -1; return -1;
} }
if (retDevCnt) { if (retDevCnt) {
// std::cout << "Total number of devices visible in system is "<< devCount << std::endl;
std::cout << devCount << std::endl; std::cout << devCount << std::endl;
} }
if (retDevInfo) { if (device != INT_MAX) {
hipDevice_t deviceT; hipDevice_t deviceT;
hipDeviceGet(&deviceT, device); hipDeviceGet(&deviceT, device);
projects/clr/hipamd/tests/src/hipEnvVarDriver.cpp
+8
View File
@@ -30,6 +30,14 @@ OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWA
#include "hip/hip_runtime.h" #include "hip/hip_runtime.h"
#include <chrono> #include <chrono>
#include <thread> #include <thread>
#ifdef _WIN64
#define popen(x,y) _popen(x,y)
#define pclose(x) _pclose(x)
#define setenv(x,y,z) _putenv_s(x,y)
#define unsetenv(x) _putenv_s(x,"")
#endif
using namespace std; using namespace std;
int getDeviceNumber() { int getDeviceNumber() {
projects/clr/hipamd/tests/src/kernel/hipLaunchParm.cpp
+4 -9
View File
@@ -577,16 +577,11 @@ __global__ void vAdd(float* a) {}
// Some wrapper macro for testing: // Some wrapper macro for testing:
#define WRAP(...) __VA_ARGS__ #define WRAP(...) __VA_ARGS__
#include <sys/time.h> #define MY_LAUNCH_MACRO(cmd, elapsed, quiet) \
#define GPU_PRINT_TIME(cmd, elapsed, quiet) \
do { \ do { \
struct timeval start, stop; \
float elapsed; \
gettimeofday(&start, NULL); \
hipDeviceSynchronize(); \ hipDeviceSynchronize(); \
cmd; \ cmd; \
hipDeviceSynchronize(); \ hipDeviceSynchronize(); \
gettimeofday(&stop, NULL); \
} while (0); } while (0);
@@ -918,14 +913,14 @@ int main() {
// Test: Passing hipLaunchKernelGGL inside another macro: // Test: Passing hipLaunchKernelGGL inside another macro:
float e0; float e0;
GPU_PRINT_TIME(hipLaunchKernelGGL(vAdd, dim3(1024), MY_LAUNCH_MACRO(hipLaunchKernelGGL(vAdd, dim3(1024),
dim3(1), 0, 0, Ad), e0, j); dim3(1), 0, 0, Ad), e0, j);
GPU_PRINT_TIME(WRAP(hipLaunchKernelGGL(vAdd, dim3(1024), MY_LAUNCH_MACRO(WRAP(hipLaunchKernelGGL(vAdd, dim3(1024),
dim3(1), 0, 0, Ad)), e0, j); dim3(1), 0, 0, Ad)), e0, j);
#ifdef EXTRA_PARENS_1 #ifdef EXTRA_PARENS_1
// Don't wrap hipLaunchKernelGGL in extra set of parens: // Don't wrap hipLaunchKernelGGL in extra set of parens:
GPU_PRINT_TIME((hipLaunchKernelGGL(vAdd, dim3(1024), MY_LAUNCH_MACRO((hipLaunchKernelGGL(vAdd, dim3(1024),
dim3(1), 0, 0, Ad)), e0, j); dim3(1), 0, 0, Ad)), e0, j);
#endif #endif
projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemcpy_simple.cpp
+3
View File
@@ -29,6 +29,9 @@ THE SOFTWARE.
#include "hip/hip_runtime.h" #include "hip/hip_runtime.h"
#include "test_common.h" #include "test_common.h"
#ifdef _WIN64
#define aligned_alloc _aligned_malloc
#endif
bool p_async = false; bool p_async = false;
projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemoryAllocateCoherentDriver.cpp
+7
View File
@@ -29,6 +29,13 @@ OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWA
#include <assert.h> #include <assert.h>
#include <string> #include <string>
#include "hip/hip_runtime.h" #include "hip/hip_runtime.h"
#ifdef _WIN64
#define popen(x,y) _popen(x,y)
#define pclose(x) _pclose(x)
#define setenv(x,y,z) _putenv_s(x,y)
#endif
using namespace std; using namespace std;
string getRes() { string getRes() {