AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

Merge pull request #1315 from emankov/clang

Просмотр исходного кода 
[HIPIFY][cuRAND][#1257] Fix
Этот коммит содержится в:
Evgeny Mankov
2019-08-09 21:35:33 +03:00
коммит произвёл GitHub
родитель eae8406e63 3ac3b2800b
Коммит ffbdb1c8af
3 изменённых файлов: 36 добавлений и 22 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
docs/markdown/CURAND_API_supported_by_HIP.md
+18 -10
Просмотреть файл
@@ -93,16 +93,24 @@
| 11 |*`CURAND_3RD`* | |
| 12 |*`CURAND_DEFINITION`* | |
| 13 |*`CURAND_POISSON`* | |
| struct | `curandStateMtgp32_t` | `hiprandStateMtgp32_t` |
| struct | `curandStateScrambledSobol64_t` | |
| struct | `curandStateSobol64_t` | |
| struct | `curandStateScrambledSobol32_t` | |
| struct | `curandStateSobol32_t` | `hiprandStateSobol32_t` |
| struct | `curandStateMRG32k3a_t` | `hiprandStateMRG32k3a_t` |
| struct | `curandStatePhilox4_32_10_t` | `hiprandStatePhilox4_32_10_t` |
| struct | `curandStateXORWOW_t` | `hiprandStateXORWOW_t` |
| struct | `curandState_t` | `hiprandState_t` |
| struct | `curandState` | `hiprandState_t` |
| struct | `curandStateMtgp32` | `hiprandStateMtgp32` |
| typedef | `curandStateMtgp32_t` | `hiprandStateMtgp32_t` |
| struct | `curandStateScrambledSobol64` | |
| typedef | `curandStateScrambledSobol64_t` | |
| struct | `curandStateSobol64` | |
| typedef | `curandStateSobol64_t` | |
| struct | `curandStateScrambledSobol32` | |
| typedef | `curandStateScrambledSobol32_t` | |
| struct | `curandStateSobol32` | `hiprandStateSobol32` |
| 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**
hipify-clang/src/CUDA2HIP_RAND_API_types.cpp
+9 -2
Просмотреть файл
@@ -53,17 +53,24 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RAND_TYPE_NAME_MAP{
{"curandDirectionVectors32_t", {"hiprandDirectionVectors32_t", "", CONV_TYPE, API_RAND}},
{"curandDirectionVectors64_t", {"hiprandDirectionVectors64_t", "", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}},
// cuRAND types for Device functions
{"curandStateMtgp32", {"hiprandStateMtgp32", "", 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}},
{"curandStateSobol64", {"hiprandStateSobol64", "", 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}},
{"curandStateSobol32", {"hiprandStateSobol32", "", 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}},
{"curandStatePhilox4_32_10", {"hiprandStatePhilox4_32_10", "", 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}},
{"curandState", {"hiprandState", "", 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)
{"CURAND_STATUS_SUCCESS", {"HIPRAND_STATUS_SUCCESS", "", CONV_NUMERIC_LITERAL, API_RAND}},
tests/hipify-clang/unit_tests/libraries/cuRAND/poisson_api_example.cu
+9 -10
Просмотреть файл
@@ -52,7 +52,7 @@ API_TYPE api;
int report_break;
int cashiers_load_h[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)
{
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;
}
// CHECK: __global__ void simple_device_API_kernel(hiprandState_t *state,
__global__ void simple_device_API_kernel(curandState *state,
// CHECK: __global__ void simple_device_API_kernel(hiprandState *state,
__global__ void simple_device_API_kernel(curandState *state,
unsigned int *queue_lengths, size_t pitch)
{
int id = threadIdx.x + blockIdx.x * blockDim.x;
unsigned int new_customers;
unsigned int queue_length = 0;
/* Copy state to local memory for efficiency */
// CHECK: hiprandState_t localState = state[id];
// CHECK: hiprandState localState = state[id];
curandState localState = state[id];
/* Simulate queue in time */
for(int min = 1; min <= 60 * HOURS; min++) {
@@ -120,7 +120,7 @@ __global__ void host_API_kernel(unsigned int *poisson_numbers,
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_2,
// CHECK: hiprandDiscreteDistribution_t poisson_3,
@@ -134,7 +134,7 @@ __global__ void robust_device_API_kernel(curandState *state,
unsigned int new_customers;
unsigned int queue_length = 0;
/* Copy state to local memory for efficiency */
// CHECK: hiprandState_t localState = state[id];
// CHECK: hiprandState localState = state[id];
curandState localState = state[id];
/* Simulate queue in time */
/* first 3 hours */
@@ -165,7 +165,7 @@ __global__ void robust_device_API_kernel(curandState *state,
curand_discrete(&localState, poisson_3);
/* Update queue */
update_queue(id, min, new_customers, queue_length,
queue_lengths, pitch);
queue_lengths, pitch);
}
/* Copy state back to global memory */
state[id] = localState;
@@ -298,12 +298,11 @@ void print_statistics(unsigned int *hostResults, size_t pitch)
}
}
int main(int argc, char *argv[])
{
int n;
size_t pitch;
// CHECK: hiprandState_t *devStates;
// CHECK: hiprandState *devStates;
curandState *devStates;
unsigned int *devResults, *hostResults;
unsigned int *poisson_numbers_d;
@@ -328,7 +327,7 @@ int main(int argc, char *argv[])
/* Allocate space for prng states on device */
// CHECK: CUDA_CALL(hipMalloc((void **)&devStates, 64 * 64 *
// CHECK: sizeof(hiprandState_t)));
// CHECK: sizeof(hiprandState)));
CUDA_CALL(cudaMalloc((void **)&devStates, 64 * 64 *
sizeof(curandState)));