[ROCm] add support for ROCm/HIP device #6086

jeffdaily · 2023-09-08T16:21:34Z

To build for ROCm:

./helpers/hipify.sh
mkdir build
cd build
cmake -DUSE_ROCM=1 ..

CUDA source files are hipified in-place using the helper script before running cmake. The "cuda" device is re-used for rocm, so device=cuda will work the same for rocm builds.

Summary of changes:

CMakeLists.txt ROCm updates, also replace glob with explicit file list
support both warpSize 32 and 64
helpers/hipify.sh script added
.gitignore to ignore generated hip source files *.prehip
disable compiler warnings
move __device__ template function PercentileDevice into header
bug fixes for __host__ __define__

- CMakeLists.txt ROCm updates, also replace glob with explicit file list - initial warpSize interop changes - helpers/hipify.sh script added - .gitignore to ignore generated hip source files

- disable compiler warnings - move PercentileDevice __device__ template function into header - bug fixes for __host__ __define__ and __HIP__ preprocessor symbols

jameslamb

Thanks for your interest in LightGBM. Since I'm not aware of any prior conversation in this project about adding support like this, we have some questions before spending time supporting this.

what is ROCm/HIP? Where can we read to learn more?
what is the value of this addition to LightGBM's users? What does this offer that the OpenCL-based and CUDA-based builds of LightGBM don't already offer?
- this project's OpenCL-based GPU build is already struggling from a severe lack of maintenance... I'm very skeptical of taking on a third GPU build
how might we test this? What types of devices should we expect to be supported?

shiyu1994 · 2023-09-08T16:52:52Z

@jeffdaily Thank you, this is very exciting! @jameslamb ROCm is the counterpart of CUDA for AMD GPU. I don't have any prior discussion with @jeffdaily about this. But it is very exciting if we can enlarge the devices supported by LightGBM.

jeffdaily · 2023-09-08T17:22:59Z

Apologies for coming out of nowhere with this. We use LightGBM; the OpenCL-based 'gpu' device already works on our AMD GPUs. But we were curious if we could get better performance if we ported the 'cuda' device to AMD GPUs. This started as a proof of concept, but it seemed useful to share even in its current state.

Using the GPU-Tutorial, here are my results on our MI210.

what is evaluated	CPU	GPU/OpenCL	"cuda" but really ROCm
correctness	auc : 0.821268 18.547533 seconds	auc : 0.821268 20.386780 seconds	auc : 0.821268 9.049307 seconds
speed objective=binary metric=auc	22.604444 seconds	18.028674 seconds	7.787303 seconds
speed objective=regression_l2 metric=l2	18.961535 seconds	14.491217 seconds	7.871302 seconds

jeffdaily · 2023-09-08T17:55:11Z

what is ROCm/HIP? Where can we read to learn more?

https://rocm.docs.amd.com/en/latest/rocm.html

what is the value of this addition to LightGBM's users? What does this offer that the OpenCL-based and CUDA-based builds of LightGBM don't already offer?

See the perf results from the comment above.

this project's OpenCL-based GPU build is already struggling from a severe lack of maintenance... I'm very skeptical of taking on a third GPU build

how might we test this? What types of devices should we expect to be supported?

Here is the current list of supported AMD GPUs.

To test this, you'll need to run on one of the supported AMD GPUs. How is the cuda device currently tested?

ibustany · 2023-09-08T19:20:35Z

Thank you and kudos Jeff!
This work has been much needed!
Best regards,
Ismail

jameslamb · 2023-09-08T20:03:09Z

To test this, you'll need to run on one of the supported AMD GPUs. How is the cuda device currently tested?

We run a VM in Azure with a Tesla V100 on it, and schedule jobs onto it via GitHub Actions.

example build link: https://github.com/microsoft/LightGBM/actions/runs/6123938185/job/16622920873#step:5:34-51
configuration:

LightGBM/.github/workflows/cuda.yml

Line 25 in 04b66e0

runs-on: [self-hosted, linux]

Are you aware of any free CI service supporting AMD GPUs? Otherwise, since I see you work for AMD and since merging this might further AMD's interests... would AMD maybe be willing to fund testing resources for this project? Maybe that's something you and @shiyu1994 (the only maintainer here who's employed by Microsoft) could coordinate?

jeffdaily · 2023-09-08T20:09:14Z

Are you aware of any free CI service supporting AMD GPUs? Otherwise, since I see you work for AMD and since merging this might further AMD's interests... would AMD maybe be willing to fund testing resources for this project? Maybe that's something you and @shiyu1994 (the only maintainer here who's employed by Microsoft) could coordinate?

Microsoft does have an AMD GPU deployment. I'm aware of it being used for onnxruntime CI purposes. I wonder if some of those resources could be used here? @shiyu1994?

jeffdaily · 2023-09-08T22:45:42Z

Noting that the only CI failure currently is not related to my changes. It seems to be a perhaps temporary environment setup issue for that job.

shiyu1994 · 2023-09-13T03:34:28Z

I have access to some AMD MI100 GPUs. But we still need separate budget for an agent with an AMD GPU if we want to test automatically in ci. Do you think it is acceptable if I run the tests for AMD GPU offline without an additional agent for ci? Given that the code for GPU version is shared by both CUDA and ROCm. @jameslamb @guolinke @jeffdaily.

jameslamb · 2023-09-13T14:10:08Z

Do you think it is acceptable if I run the tests for AMD GPU offline without an additional agent for ci?

If you feel confident in these changes based on that, and you think the added complexity in the CUDA code is worth it, that's fine with me. I'll defer to your opinion.

But without a CI job, there's a high risk that future refactorings will break this support again.

dismissing

jameslamb · 2023-09-13T14:12:53Z

I dismissed my review, so that it doesn't block merging. My initial questions have been answered, thanks very much for those links and all that information!

@shiyu1994 and @guolinke seem excited about this addition... that's good enough for me 😊

I'll defer to them to review the code, as I know very little about CUDA.

shiyu1994 · 2023-10-08T15:26:20Z

@jeffdaily Thanks for the great work! I'll review this in the next few days.

StrikerRUS

Thanks for updating the code!

I think this PR is blocked by #6766.

Also, I searched the CUDA code in the repo for the literals 32 and 64 and left some comments in places where warpsize can potentially be adjusted.

CMakeLists.txt

include/LightGBM/cuda/cuda_rocm_interop.h

src/treelearner/cuda/cuda_histogram_constructor.cu

StrikerRUS · 2024-12-24T21:27:52Z

src/treelearner/cuda/cuda_single_gpu_tree_learner.cu

@@ -167,7 +169,7 @@ void CUDASingleGPUTreeLearner::LaunchReduceLeafStatKernel(

 template <typename T, bool IS_INNER>
 __global__ void CalcBitsetLenKernel(const CUDASplitInfo* best_split_info, size_t* out_len_buffer) {
-  __shared__ size_t shared_mem_buffer[32];
+  __shared__ size_t shared_mem_buffer[WARPSIZE];


Should we also adjust the code that is relying on warpsize is always 32? For example, here:

LightGBM/src/treelearner/cuda/cuda_single_gpu_tree_learner.cu

Lines 181 to 183 in 60b0155

len = (val / 32) + 1;

}

const size_t block_max_len = ShuffleReduceMax<size_t>(len, shared_mem_buffer, blockDim.x);

Done in 3cd34a2.

Could you please explain how exactly this my comment was addressed in 3cd34a2?

src/treelearner/cuda/cuda_data_partition.cu

StrikerRUS · 2024-12-24T21:57:01Z

src/objective/cuda/cuda_rank_objective.cu

@@ -354,7 +358,7 @@ void CUDALambdarankNDCG::LaunchGetGradientsKernel(const double* score, score_t*
    }
  } else {
    BitonicArgSortItemsGlobal(score, num_queries_, cuda_query_boundaries_, cuda_item_indices_buffer_.RawData());
-    if (num_rank_label <= 32) {
+    if (num_rank_label <= 32 && device_prop.warpSize == 32) {


Should we adjust the following code for warpsize other than 32?

LightGBM/src/objective/cuda/cuda_rank_objective.cu

Lines 407 to 408 in 60b0155

// assert that warpSize == 32

__shared__ double shared_buffer[32];

LightGBM/src/objective/cuda/cuda_rank_objective.cu

Lines 525 to 526 in 60b0155

// assert that warpSize == 32, so we use buffer size 1024 / 32 = 32

__shared__ double shared_buffer[32];

Done in 3cd34a2.

I guess that according to the initial comment it should be shared_buffer[1024 / WARPSIZE] but not shared_buffer[WARPSIZE]

Co-authored-by: Nikita Titov <[email protected]>

use WARPSIZE

shiyu1994 · 2025-03-04T17:30:31Z

@StrikerRUS @jameslamb @guolinke Could you help to review this again when you have time? Thanks.

StrikerRUS

@shiyu1994 Thank you for considering my comments! I think that two of them were not addressed by your recent refactoring. Please check #6086 (comment) and #6086 (comment).

shiyu1994 · 2025-06-10T03:45:43Z

@shiyu1994 Thank you for considering my comments! I think that two of them were not addressed by your recent refactoring. Please check #6086 (comment) and #6086 (comment).

Thank you for the very careful check. I've done the fixes in 28d4648. Could you please review it again? @StrikerRUS

StrikerRUS

@shiyu1994

I've done the fixes in 28d4648.

I believe that right commit would be 6732b79.

I left my last suggestion for this PR below.

StrikerRUS · 2025-06-15T19:51:30Z

src/objective/cuda/cuda_rank_objective.cu

@@ -405,7 +409,7 @@ __global__ void GetGradientsKernel_RankXENDCG_SharedMemory(
    const data_size_t block_reduce_size = query_item_count >= 1024 ? 1024 : query_item_count;
    __shared__ double shared_rho[SHARED_MEMORY_SIZE];
    // assert that warpSize == 32
-    __shared__ double shared_buffer[32];
+    __shared__ double shared_buffer[WARPSIZE];


I guess it should be 1024 / WARPSIZE similarly to L530 in this file.

Done via d4676d9.
Please check.
Thanks for the finding!

shiyu1994 · 2025-06-17T06:57:03Z

@StrikerRUS Thanks for your review. Could you please check this PR again? If there's no other problems, let's merge this.

shiyu1994 · 2025-06-17T09:16:10Z

/AzurePipelines run

StrikerRUS

@jeffdaily Thank you so much for proposing this PR!

And thanks a lot @shiyu1994 for finishing it!

StrikerRUS · 2025-06-22T19:20:46Z

@jameslamb Could you please refresh your blocking review?

#6086 (review)

jeffdaily added 4 commits September 7, 2023 20:48

[ROCm] add support for ROCm/HIP

01ff268

- CMakeLists.txt ROCm updates, also replace glob with explicit file list - initial warpSize interop changes - helpers/hipify.sh script added - .gitignore to ignore generated hip source files

more rocm updates

ac966a5

- disable compiler warnings - move PercentileDevice __device__ template function into header - bug fixes for __host__ __define__ and __HIP__ preprocessor symbols

more bug fixes

6ae6432

warp 32 vs 64 updates

5b87bcd

jeffdaily requested review from guolinke, jameslamb, shiyu1994 and jmoralez as code owners September 8, 2023 16:21

jameslamb previously requested changes Sep 8, 2023

View reviewed changes

jameslamb added in progress feature labels Sep 8, 2023

lint fixes

3ad89f8

missing device_index variable

7b8b6a0

accidental inclusion of hip headers

62aa30b

jeffdaily and others added 3 commits September 11, 2023 16:21

copyright notice compliance

bb27c55

Merge branch 'master' into rocm2

58ace9c

Merge branch 'master' into rocm2

0bc1cfb

Merge branch 'master' into rocm2

a7c9653

shiyu1994 added 2 commits December 24, 2024 07:44

Merge branch 'rocm2' of https://github.com/jeffdaily/LightGBM into HEAD

47fc353

Merge branch 'master' into rocm2

2e8869c

shiyu1994 requested review from StrikerRUS and jameslamb December 24, 2024 07:46

StrikerRUS reviewed Dec 24, 2024

View reviewed changes

shiyu1994 and others added 4 commits February 5, 2025 17:25

Update CMakeLists.txt

b173124

Co-authored-by: Nikita Titov <[email protected]>

Update CMakeLists.txt

785b341

Co-authored-by: Nikita Titov <[email protected]>

merge master

f4c605d

clean up CMakeLists.txt

3cd34a2

use WARPSIZE

shiyu1994 requested a review from StrikerRUS February 26, 2025 03:12

Merge branch 'master' into rocm2

89605b8

jeffdaily dismissed shiyu1994’s stale review via 89605b8 April 1, 2025 00:33

StrikerRUS requested changes Apr 13, 2025

View reviewed changes

shiyu1994 and others added 2 commits April 16, 2025 12:40

Merge branch 'master' into rocm2

c66b8f3

Merge branch 'master' into rocm2

8591248

StrikerRUS mentioned this pull request Apr 20, 2025

[CUDA] Multi-GPU for CUDA Version #6138

Open

shiyu1994 and others added 3 commits June 10, 2025 11:38

Merge branch 'master' into rocm2

8b1c95a

use WARPSIZE

6732b79

Merge branch 'rocm2' of https://github.com/jeffdaily/LightGBM into HEAD

28d4648

StrikerRUS reviewed Jun 15, 2025

View reviewed changes

shiyu1994 added 2 commits June 17, 2025 14:42

Merge branch 'master' into rocm2

9a19de6

fix share buffer size

d4676d9

shiyu1994 added a commit that referenced this pull request Jun 17, 2025

remove WARPSIZE before #6086 is merged

c29082e

StrikerRUS removed the awaiting review label Jun 22, 2025

StrikerRUS approved these changes Jun 22, 2025

View reviewed changes

	len = (val / 32) + 1;
	}
	const size_t block_max_len = ShuffleReduceMax<size_t>(len, shared_mem_buffer, blockDim.x);

	// assert that warpSize == 32
	__shared__ double shared_buffer[32];

	// assert that warpSize == 32, so we use buffer size 1024 / 32 = 32
	__shared__ double shared_buffer[32];

[ROCm] add support for ROCm/HIP device #6086

Are you sure you want to change the base?

[ROCm] add support for ROCm/HIP device #6086

Uh oh!

Conversation

jeffdaily commented Sep 8, 2023

Uh oh!

jameslamb left a comment

Choose a reason for hiding this comment

Uh oh!

shiyu1994 commented Sep 8, 2023

Uh oh!

jeffdaily commented Sep 8, 2023

Uh oh!

jeffdaily commented Sep 8, 2023

Uh oh!

ibustany commented Sep 8, 2023

Uh oh!

jameslamb commented Sep 8, 2023

Uh oh!

jeffdaily commented Sep 8, 2023

Uh oh!

jeffdaily commented Sep 8, 2023

Uh oh!

shiyu1994 commented Sep 13, 2023

Uh oh!

jameslamb commented Sep 13, 2023

Uh oh!

jameslamb commented Sep 13, 2023

Uh oh!

shiyu1994 commented Oct 8, 2023

Uh oh!

StrikerRUS left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

shiyu1994 commented Mar 4, 2025

Uh oh!

StrikerRUS left a comment

Choose a reason for hiding this comment

Uh oh!

shiyu1994 commented Jun 10, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

StrikerRUS left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

shiyu1994 commented Jun 17, 2025

Uh oh!

shiyu1994 commented Jun 17, 2025

Uh oh!

StrikerRUS left a comment

Choose a reason for hiding this comment

Uh oh!

StrikerRUS commented Jun 22, 2025

Uh oh!

Uh oh!

shiyu1994 commented Jun 10, 2025 •

edited

Loading