Skip to content
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.

Sign up for GitHub

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

Jump to bottom

[ROCm] Adding ROCM support for the population_count op #29602

Merged
merged 1 commit into from Jun 12, 2019
Merged

[ROCm] Adding ROCM support for the population_count op #29602

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Jump to
Jump to file
Failed to load files.
Diff view
Diff view
4 changes: 2 additions & 2 deletions tensorflow/core/kernels/population_count_op.cc
View file
Open in desktop
Expand Up @@ -122,7 +122,7 @@ struct PopulationCount<CPUDevice, T> {

} // namespace functor

#if GOOGLE_CUDA
#if GOOGLE_CUDA || TENSORFLOW_USE_ROCM

#define REGISTER_POPULATION_COUNT(type) \
REGISTER_KERNEL_BUILDER( \
Expand Down Expand Up @@ -158,6 +158,6 @@ TF_CALL_int64(DECLARE_GPU_SPEC);

} // namespace functor

#endif // GOOGLE_CUDA
#endif // GOOGLE_CUDA || TENSORFLOW_USE_ROCM

} // namespace tensorflow
16 changes: 8 additions & 8 deletions tensorflow/core/kernels/population_count_op_gpu.cu.cc
View file
Open in desktop
Expand Up @@ -14,7 +14,7 @@ See the License for the specific language governing permissions and
limitations under the License.
==============================================================================*/

#if GOOGLE_CUDA
#if GOOGLE_CUDA || TENSORFLOW_USE_ROCM

#define EIGEN_USE_GPU

Expand All @@ -35,14 +35,14 @@ namespace functor {
template <typename T>
__global__ void PopulationCountKernel(const int size, const T* input,
uint8* output) {
CUDA_1D_KERNEL_LOOP(i, size) { output[i] = __popc(ldg(input + i)); }
GPU_1D_KERNEL_LOOP(i, size) { output[i] = __popc(ldg(input + i)); }
}

template <>
__global__ void PopulationCountKernel(const int size, const int8* input,
uint8* output) {
// For some reason, __popc on a negative int8 gets confused.
CUDA_1D_KERNEL_LOOP(i, size) {
GPU_1D_KERNEL_LOOP(i, size) {
output[i] = __popc(ldg(reinterpret_cast<const uint8*>(input + i)));
}
}
Expand All @@ -51,15 +51,15 @@ template <>
__global__ void PopulationCountKernel(const int size, const int16* input,
uint8* output) {
// For some reason, __popc on a negative int16 gets confused.
CUDA_1D_KERNEL_LOOP(i, size) {
GPU_1D_KERNEL_LOOP(i, size) {
output[i] = __popc(ldg(reinterpret_cast<const uint16*>(input + i)));
}
}

template <>
__global__ void PopulationCountKernel<int64>(const int size, const int64* input,
uint8* output) {
CUDA_1D_KERNEL_LOOP(i, size) { output[i] = __popcll(ldg(input + i)); }
GPU_1D_KERNEL_LOOP(i, size) { output[i] = __popcll(ldg(input + i)); }
}

#define DEFINE_GPU_SPECS(T) \
Expand All @@ -69,8 +69,8 @@ __global__ void PopulationCountKernel<int64>(const int size, const int64* input,
TTypes<uint8>::Flat output) { \
const GPUDevice& d = c->eigen_device<GPUDevice>(); \
int64 total_count = input.size(); \
GpuLaunchConfig config = GetCudaLaunchConfig(total_count, d); \
TF_CHECK_OK(CudaLaunchKernel(PopulationCountKernel<T>, config.block_count, \
GpuLaunchConfig config = GetGpuLaunchConfig(total_count, d); \
TF_CHECK_OK(GpuLaunchKernel(PopulationCountKernel<T>, config.block_count, \
config.thread_per_block, 0, d.stream(), \
total_count, input.data(), output.data())); \
}
Expand All @@ -88,4 +88,4 @@ TF_CALL_int64(DEFINE_GPU_SPECS);

} // namespace tensorflow

#endif // GOOGLE_CUDA
#endif // GOOGLE_CUDA || TENSORFLOW_USE_ROCM