diff --git a/docs/markdown/CURAND_API_supported_by_HIP.md b/docs/markdown/CURAND_API_supported_by_HIP.md index 0561dd6144..83c41e59ff 100644 --- a/docs/markdown/CURAND_API_supported_by_HIP.md +++ b/docs/markdown/CURAND_API_supported_by_HIP.md @@ -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** diff --git a/hipify-clang/src/CUDA2HIP_RAND_API_types.cpp b/hipify-clang/src/CUDA2HIP_RAND_API_types.cpp index 6bf3223cb0..e6eae1db69 100644 --- a/hipify-clang/src/CUDA2HIP_RAND_API_types.cpp +++ b/hipify-clang/src/CUDA2HIP_RAND_API_types.cpp @@ -53,17 +53,24 @@ const std::map 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}}, diff --git a/tests/hipify-clang/unit_tests/libraries/cuRAND/poisson_api_example.cu b/tests/hipify-clang/unit_tests/libraries/cuRAND/poisson_api_example.cu index 4e748973d5..f4fd05ba48 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuRAND/poisson_api_example.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuRAND/poisson_api_example.cu @@ -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)));