HIP icon indicating copy to clipboard operation
HIP copied to clipboard

[Issue]: hiprtcCompileProgram() fails to compile a kernel that includes hiprand code

Open macos-fuse-t opened this issue 7 months ago • 2 comments

Problem Description

I'm trying to dynamically compile the following kernel containing calls to the hiprand library.

#include <hiprand/hiprand_kernel.h>

const int N = 10240;           // total number of randoms
const int THREADS_PER_BLOCK = 256;

__global__ void generate_random_kernel(float *data, unsigned long seed) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx >= N) return;

    hiprandState state;
    hiprand_init(
        seed,
        idx,
        0,
        &state
    );

    data[idx] = hiprand_uniform(&state);
}

I'm getting the following compilation errors:

Compilation failed:
In file included from /tmp/comgr-a98195/input/kernel.hip:3:
In file included from /opt/rocm-6.4.0/include/hiprand/hiprand_kernel.h:29:
In file included from /opt/rocm-6.4.0/include/hiprand/hiprand.h:78:
In file included from /opt/rocm-6.4.0/include/hiprand/hiprand_rocm.h:24:
/opt/rocm-6.4.0/include/rocrand/rocrand.h:582:49: error: unknown type name 'hipStream_t'
  582 | rocrand_set_stream(rocrand_generator generator, hipStream_t stream);
      |                                                 ^
In file included from /tmp/comgr-a98195/input/kernel.hip:3:
In file included from /opt/rocm-6.4.0/include/hiprand/hiprand_kernel.h:29:
/opt/rocm-6.4.0/include/hiprand/hiprand.h:643:48: error: unknown type name 'hipStream_t'
  643 | hiprandSetStream(hiprandGenerator_t generator, hipStream_t stream);
      |                                                ^
In file included from /tmp/comgr-a98195/input/kernel.hip:3:
In file included from /opt/rocm-6.4.0/include/hiprand/hiprand_kernel.h:110:
In file included from /opt/rocm-6.4.0/include/hiprand/hiprand_kernel_rocm.h:32:
/usr/lib/gcc/x86_64-linux-gnu/12/../../../../include/c++/12/type_traits:62:12: error: redefinition of 'integral_constant'
   62 |     struct integral_constant
      |            ^
....

My test compilation code and the kernel attached

compile.txt kernel.txt

Operating System

Ubuntu 22

CPU

Intel(R) Xeon(R) Platinum 8470

GPU

MI300X

ROCm Version

6.4.0

ROCm Component

HIP

Steps to Reproduce

compile.txt kernel.txt

(Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support

compile.txt kernel.txt Attached the compilation code and the kernel

Additional Information

No response

macos-fuse-t avatar May 20 '25 08:05 macos-fuse-t

Compiling of rocm libs inside hiprtc is not supported. We do not test for it. But let me see if there is something I can do here.

cjatin avatar May 20 '25 11:05 cjatin

Workaround would be appreciated. The similar code works on the cuda stack with curand kernels and everything and I need to make it work on the rocm stack somehow as well.

macos-fuse-t avatar May 20 '25 14:05 macos-fuse-t

This issue has been migrated to: https://github.com/ROCm/rocm-systems/issues/400

systems-assistant[bot] avatar Aug 18 '25 18:08 systems-assistant[bot]