Skip to content

[CUDA][HIP] add test for uses of std::array on device code #244

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Draft
wants to merge 2 commits into
base: main
Choose a base branch
from

Conversation

jmmartinez
Copy link

First attempt at testing llvm/llvm-project#136133 .

Draft for the moment, I have to manually test this further.

__global__ void kernel() { do_test(); }

int main() {
kernel<<<32, 32>>>();
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

can std::array be parameter type of kernel or device functions?

can it be used as non-constexpr stack variable in kernel or device functions?

if so, can we have testcases for those?


#define cudaError_t hipError_t
#define cudaSuccess hipSuccess
#define cudaDeviceSynchronize hipDeviceSynchronize
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I am wondering whether we have a systematic way to use CUDA tests for HIP. Using macro to map CUDA entries to HIP is not only tedious but may impose restrictions on how those CUDA tests are written.

I am wondering whether we could hipify the .cu file instead. Currently /opt/rocm/bin/hipify-perl is available on the HIP buildbot. For each .cu file to be hipified, you can use add_custom_command to hipify it to a .hip file then test it.

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Can we try using hipify-clang instead? hipify-perl is going to be deprecated soon more than likely

@emankov

Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I've just built the upstream buildbot image from https://github.com/ROCm/aomp/tree/aomp-dev/upstream-buildbots/Ubu22 and there is no hipify in Rocm's installation (it's quite a minimal install).

I'll check if I install them if the tests work. I agree that it would be much better.

Copy link
Author

@jmmartinez jmmartinez May 15, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I've tried with hipify-perl and it works.

hipify-clang did not, it yield errors for every test.

One common error case was with multiple main definitions (even if they are on different ifdef branches):

#include <stdio.h>
#if __cplusplus < 201103L
int main() { puts("foo\n"); return 0; }
#else
int main() { puts("bar\n"); return 0; }
#endif
/tmp/foo.cu-5f2589.hip:5:5: error: redefinition of 'main'
    5 | int main() { puts("bar\n"); return 0; }
      |     ^
/tmp/foo.cu-5f2589.hip:3:5: note: previous definition is here
    3 | int main() { puts("foo\n"); return 0; }

I'll adapt these tests to run with hipify and remove the hacky header.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

3 participants