Skip to content

GitLab

".github/git@developer.sourcefind.cn:tianlh/lightgbm-dcu.git" did not exist on "2c6b74afb438f5db7c7fcf383bc14364e950bcb2"
Unverified Commit 5c9e61d1 authored by shiyu1994's avatar shiyu1994 Committed by GitHub
Browse files

[CUDA] Set GPU device ID in threads (#6028) 



* set gpu device id in open mp threads

* move SetCUDADevice outside for loop

---------
Co-authored-by: default avatarJames Lamb <jaylamb20@gmail.com>
parent fe838d88
Hide whitespace changes
Inline Side-by-side
Showing with 95 additions and 71 deletions
include/LightGBM/cuda/cuda_column_data.hpp
View file @ 5c9e61d1
...@@ -98,6 +98,7 @@ class CUDAColumnData { ...@@ -98,6 +98,7 @@ class CUDAColumnData {
void ResizeWhenCopySubrow(const data_size_t num_used_indices); void ResizeWhenCopySubrow(const data_size_t num_used_indices);
int gpu_device_id_;
int num_threads_; int num_threads_;
data_size_t num_data_; data_size_t num_data_;
int num_columns_; int num_columns_;
......
include/LightGBM/cuda/cuda_metric.hpp
View file @ 5c9e61d1
...@@ -9,6 +9,7 @@ ...@@ -9,6 +9,7 @@
#ifdef USE_CUDA #ifdef USE_CUDA
#include <LightGBM/cuda/cuda_utils.h>
#include <LightGBM/metric.h> #include <LightGBM/metric.h>
namespace LightGBM { namespace LightGBM {
...@@ -19,6 +20,8 @@ class CUDAMetricInterface: public HOST_METRIC { ...@@ -19,6 +20,8 @@ class CUDAMetricInterface: public HOST_METRIC {
explicit CUDAMetricInterface(const Config& config): HOST_METRIC(config) { explicit CUDAMetricInterface(const Config& config): HOST_METRIC(config) {
cuda_labels_ = nullptr; cuda_labels_ = nullptr;
cuda_weights_ = nullptr; cuda_weights_ = nullptr;
const int gpu_device_id = config.gpu_device_id >= 0 ? config.gpu_device_id : 0;
SetCUDADevice(gpu_device_id, __FILE__, __LINE__);
} }
void Init(const Metadata& metadata, data_size_t num_data) override { void Init(const Metadata& metadata, data_size_t num_data) override {
......
include/LightGBM/cuda/cuda_objective_function.hpp
View file @ 5c9e61d1
...@@ -21,7 +21,10 @@ namespace LightGBM { ...@@ -21,7 +21,10 @@ namespace LightGBM {
template <typename HOST_OBJECTIVE> template <typename HOST_OBJECTIVE>
class CUDAObjectiveInterface: public HOST_OBJECTIVE { class CUDAObjectiveInterface: public HOST_OBJECTIVE {
public: public:
explicit CUDAObjectiveInterface(const Config& config): HOST_OBJECTIVE(config) {} explicit CUDAObjectiveInterface(const Config& config): HOST_OBJECTIVE(config) {
const int gpu_device_id = config.gpu_device_id >= 0 ? config.gpu_device_id : 0;
SetCUDADevice(gpu_device_id, __FILE__, __LINE__);
}
explicit CUDAObjectiveInterface(const std::vector<std::string>& strs): HOST_OBJECTIVE(strs) {} explicit CUDAObjectiveInterface(const std::vector<std::string>& strs): HOST_OBJECTIVE(strs) {}
......
include/LightGBM/cuda/cuda_utils.h
View file @ 5c9e61d1
...@@ -28,6 +28,8 @@ inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = ...@@ -28,6 +28,8 @@ inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort =
void SetCUDADevice(int gpu_device_id, const char* file, int line); void SetCUDADevice(int gpu_device_id, const char* file, int line);
int GetCUDADevice(const char* file, int line);
template <typename T> template <typename T>
void AllocateCUDAMemory(T** out_ptr, size_t size, const char* file, const int line) { void AllocateCUDAMemory(T** out_ptr, size_t size, const char* file, const int line) {
void* tmp_ptr = nullptr; void* tmp_ptr = nullptr;
......
src/cuda/cuda_utils.cpp
View file @ 5c9e61d1
...@@ -26,6 +26,12 @@ void SetCUDADevice(int gpu_device_id, const char* file, int line) { ...@@ -26,6 +26,12 @@ void SetCUDADevice(int gpu_device_id, const char* file, int line) {
} }
} }
int GetCUDADevice(const char* file, int line) {
int cur_gpu_device_id = 0;
CUDASUCCESS_OR_FATAL_OUTER(cudaGetDevice(&cur_gpu_device_id));
return cur_gpu_device_id;
}
} // namespace LightGBM } // namespace LightGBM
#endif // USE_CUDA #endif // USE_CUDA
src/io/cuda/cuda_column_data.cpp
View file @ 5c9e61d1
...@@ -12,11 +12,8 @@ namespace LightGBM { ...@@ -12,11 +12,8 @@ namespace LightGBM {
CUDAColumnData::CUDAColumnData(const data_size_t num_data, const int gpu_device_id) { CUDAColumnData::CUDAColumnData(const data_size_t num_data, const int gpu_device_id) {
num_threads_ = OMP_NUM_THREADS(); num_threads_ = OMP_NUM_THREADS();
num_data_ = num_data; num_data_ = num_data;
if (gpu_device_id >= 0) { gpu_device_id_ = gpu_device_id >= 0 ? gpu_device_id : 0;
SetCUDADevice(gpu_device_id, __FILE__, __LINE__); SetCUDADevice(gpu_device_id_, __FILE__, __LINE__);
} else {
SetCUDADevice(0, __FILE__, __LINE__);
}
cuda_used_indices_ = nullptr; cuda_used_indices_ = nullptr;
cuda_data_by_column_ = nullptr; cuda_data_by_column_ = nullptr;
cuda_column_bit_type_ = nullptr; cuda_column_bit_type_ = nullptr;
...@@ -117,37 +114,41 @@ void CUDAColumnData::Init(const int num_columns, ...@@ -117,37 +114,41 @@ void CUDAColumnData::Init(const int num_columns,
feature_mfb_is_na_ = feature_mfb_is_na; feature_mfb_is_na_ = feature_mfb_is_na;
data_by_column_.resize(num_columns_, nullptr); data_by_column_.resize(num_columns_, nullptr);
OMP_INIT_EX(); OMP_INIT_EX();
#pragma omp parallel for schedule(static) num_threads(num_threads_) #pragma omp parallel num_threads(num_threads_)
for (int column_index = 0; column_index < num_columns_; ++column_index) { {
OMP_LOOP_EX_BEGIN(); SetCUDADevice(gpu_device_id_, __FILE__, __LINE__);
const int8_t bit_type = column_bit_type[column_index]; #pragma omp for schedule(static)
if (column_data[column_index] != nullptr) { for (int column_index = 0; column_index < num_columns_; ++column_index) {
// is dense column OMP_LOOP_EX_BEGIN();
if (bit_type == 4) { const int8_t bit_type = column_bit_type[column_index];
column_bit_type_[column_index] = 8; if (column_data[column_index] != nullptr) {
InitOneColumnData<false, true, uint8_t>(column_data[column_index], nullptr, &data_by_column_[column_index]); // is dense column
} else if (bit_type == 8) { if (bit_type == 4) {
InitOneColumnData<false, false, uint8_t>(column_data[column_index], nullptr, &data_by_column_[column_index]); column_bit_type_[column_index] = 8;
} else if (bit_type == 16) { InitOneColumnData<false, true, uint8_t>(column_data[column_index], nullptr, &data_by_column_[column_index]);
InitOneColumnData<false, false, uint16_t>(column_data[column_index], nullptr, &data_by_column_[column_index]); } else if (bit_type == 8) {
} else if (bit_type == 32) { InitOneColumnData<false, false, uint8_t>(column_data[column_index], nullptr, &data_by_column_[column_index]);
InitOneColumnData<false, false, uint32_t>(column_data[column_index], nullptr, &data_by_column_[column_index]); } else if (bit_type == 16) {
} else { InitOneColumnData<false, false, uint16_t>(column_data[column_index], nullptr, &data_by_column_[column_index]);
Log::Fatal("Unknow column bit type %d", bit_type); } else if (bit_type == 32) {
} InitOneColumnData<false, false, uint32_t>(column_data[column_index], nullptr, &data_by_column_[column_index]);
} else { } else {
// is sparse column Log::Fatal("Unknow column bit type %d", bit_type);
if (bit_type == 8) { }
InitOneColumnData<true, false, uint8_t>(nullptr, column_bin_iterator[column_index], &data_by_column_[column_index]);
} else if (bit_type == 16) {
InitOneColumnData<true, false, uint16_t>(nullptr, column_bin_iterator[column_index], &data_by_column_[column_index]);
} else if (bit_type == 32) {
InitOneColumnData<true, false, uint32_t>(nullptr, column_bin_iterator[column_index], &data_by_column_[column_index]);
} else { } else {
Log::Fatal("Unknow column bit type %d", bit_type); // is sparse column
if (bit_type == 8) {
InitOneColumnData<true, false, uint8_t>(nullptr, column_bin_iterator[column_index], &data_by_column_[column_index]);
} else if (bit_type == 16) {
InitOneColumnData<true, false, uint16_t>(nullptr, column_bin_iterator[column_index], &data_by_column_[column_index]);
} else if (bit_type == 32) {
InitOneColumnData<true, false, uint32_t>(nullptr, column_bin_iterator[column_index], &data_by_column_[column_index]);
} else {
Log::Fatal("Unknow column bit type %d", bit_type);
}
} }
OMP_LOOP_EX_END();
} }
OMP_LOOP_EX_END();
} }
OMP_THROW_EX(); OMP_THROW_EX();
feature_to_column_ = feature_to_column; feature_to_column_ = feature_to_column;
...@@ -182,24 +183,28 @@ void CUDAColumnData::CopySubrow( ...@@ -182,24 +183,28 @@ void CUDAColumnData::CopySubrow(
AllocateCUDAMemory<data_size_t>(&cuda_used_indices_, num_used_indices_size, __FILE__, __LINE__); AllocateCUDAMemory<data_size_t>(&cuda_used_indices_, num_used_indices_size, __FILE__, __LINE__);
data_by_column_.resize(num_columns_, nullptr); data_by_column_.resize(num_columns_, nullptr);
OMP_INIT_EX(); OMP_INIT_EX();
#pragma omp parallel for schedule(static) num_threads(num_threads_) #pragma omp parallel num_threads(num_threads_)
for (int column_index = 0; column_index < num_columns_; ++column_index) { {
OMP_LOOP_EX_BEGIN(); SetCUDADevice(gpu_device_id_, __FILE__, __LINE__);
const uint8_t bit_type = column_bit_type_[column_index]; #pragma omp for schedule(static)
if (bit_type == 8) { for (int column_index = 0; column_index < num_columns_; ++column_index) {
uint8_t* column_data = nullptr; OMP_LOOP_EX_BEGIN();
AllocateCUDAMemory<uint8_t>(&column_data, num_used_indices_size, __FILE__, __LINE__); const uint8_t bit_type = column_bit_type_[column_index];
data_by_column_[column_index] = reinterpret_cast<void*>(column_data); if (bit_type == 8) {
} else if (bit_type == 16) { uint8_t* column_data = nullptr;
uint16_t* column_data = nullptr; AllocateCUDAMemory<uint8_t>(&column_data, num_used_indices_size, __FILE__, __LINE__);
AllocateCUDAMemory<uint16_t>(&column_data, num_used_indices_size, __FILE__, __LINE__); data_by_column_[column_index] = reinterpret_cast<void*>(column_data);
data_by_column_[column_index] = reinterpret_cast<void*>(column_data); } else if (bit_type == 16) {
} else if (bit_type == 32) { uint16_t* column_data = nullptr;
uint32_t* column_data = nullptr; AllocateCUDAMemory<uint16_t>(&column_data, num_used_indices_size, __FILE__, __LINE__);
AllocateCUDAMemory<uint32_t>(&column_data, num_used_indices_size, __FILE__, __LINE__); data_by_column_[column_index] = reinterpret_cast<void*>(column_data);
data_by_column_[column_index] = reinterpret_cast<void*>(column_data); } else if (bit_type == 32) {
uint32_t* column_data = nullptr;
AllocateCUDAMemory<uint32_t>(&column_data, num_used_indices_size, __FILE__, __LINE__);
data_by_column_[column_index] = reinterpret_cast<void*>(column_data);
}
OMP_LOOP_EX_END();
} }
OMP_LOOP_EX_END();
} }
OMP_THROW_EX(); OMP_THROW_EX();
InitCUDAMemoryFromHostMemory<void*>(&cuda_data_by_column_, data_by_column_.data(), data_by_column_.size(), __FILE__, __LINE__); InitCUDAMemoryFromHostMemory<void*>(&cuda_data_by_column_, data_by_column_.data(), data_by_column_.size(), __FILE__, __LINE__);
...@@ -221,27 +226,31 @@ void CUDAColumnData::ResizeWhenCopySubrow(const data_size_t num_used_indices) { ...@@ -221,27 +226,31 @@ void CUDAColumnData::ResizeWhenCopySubrow(const data_size_t num_used_indices) {
DeallocateCUDAMemory<data_size_t>(&cuda_used_indices_, __FILE__, __LINE__); DeallocateCUDAMemory<data_size_t>(&cuda_used_indices_, __FILE__, __LINE__);
AllocateCUDAMemory<data_size_t>(&cuda_used_indices_, num_used_indices_size, __FILE__, __LINE__); AllocateCUDAMemory<data_size_t>(&cuda_used_indices_, num_used_indices_size, __FILE__, __LINE__);
OMP_INIT_EX(); OMP_INIT_EX();
#pragma omp parallel for schedule(static) num_threads(num_threads_) #pragma omp parallel num_threads(num_threads_)
for (int column_index = 0; column_index < num_columns_; ++column_index) { {
OMP_LOOP_EX_BEGIN(); SetCUDADevice(gpu_device_id_, __FILE__, __LINE__);
const uint8_t bit_type = column_bit_type_[column_index]; #pragma omp for schedule(static)
if (bit_type == 8) { for (int column_index = 0; column_index < num_columns_; ++column_index) {
uint8_t* column_data = reinterpret_cast<uint8_t*>(data_by_column_[column_index]); OMP_LOOP_EX_BEGIN();
DeallocateCUDAMemory<uint8_t>(&column_data, __FILE__, __LINE__); const uint8_t bit_type = column_bit_type_[column_index];
AllocateCUDAMemory<uint8_t>(&column_data, num_used_indices_size, __FILE__, __LINE__); if (bit_type == 8) {
data_by_column_[column_index] = reinterpret_cast<void*>(column_data); uint8_t* column_data = reinterpret_cast<uint8_t*>(data_by_column_[column_index]);
} else if (bit_type == 16) { DeallocateCUDAMemory<uint8_t>(&column_data, __FILE__, __LINE__);
uint16_t* column_data = reinterpret_cast<uint16_t*>(data_by_column_[column_index]); AllocateCUDAMemory<uint8_t>(&column_data, num_used_indices_size, __FILE__, __LINE__);
DeallocateCUDAMemory<uint16_t>(&column_data, __FILE__, __LINE__); data_by_column_[column_index] = reinterpret_cast<void*>(column_data);
AllocateCUDAMemory<uint16_t>(&column_data, num_used_indices_size, __FILE__, __LINE__); } else if (bit_type == 16) {
data_by_column_[column_index] = reinterpret_cast<void*>(column_data); uint16_t* column_data = reinterpret_cast<uint16_t*>(data_by_column_[column_index]);
} else if (bit_type == 32) { DeallocateCUDAMemory<uint16_t>(&column_data, __FILE__, __LINE__);
uint32_t* column_data = reinterpret_cast<uint32_t*>(data_by_column_[column_index]); AllocateCUDAMemory<uint16_t>(&column_data, num_used_indices_size, __FILE__, __LINE__);
DeallocateCUDAMemory<uint32_t>(&column_data, __FILE__, __LINE__); data_by_column_[column_index] = reinterpret_cast<void*>(column_data);
AllocateCUDAMemory<uint32_t>(&column_data, num_used_indices_size, __FILE__, __LINE__); } else if (bit_type == 32) {
data_by_column_[column_index] = reinterpret_cast<void*>(column_data); uint32_t* column_data = reinterpret_cast<uint32_t*>(data_by_column_[column_index]);
DeallocateCUDAMemory<uint32_t>(&column_data, __FILE__, __LINE__);
AllocateCUDAMemory<uint32_t>(&column_data, num_used_indices_size, __FILE__, __LINE__);
data_by_column_[column_index] = reinterpret_cast<void*>(column_data);
}
OMP_LOOP_EX_END();
} }
OMP_LOOP_EX_END();
} }
OMP_THROW_EX(); OMP_THROW_EX();
DeallocateCUDAMemory<void*>(&cuda_data_by_column_, __FILE__, __LINE__); DeallocateCUDAMemory<void*>(&cuda_data_by_column_, __FILE__, __LINE__);
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment