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

[sdk] removed limits.cuh in favor of libcu++'s limits #654

Draft
wants to merge 1 commit into
base: dev
Choose a base branch
from
Draft

[sdk] removed limits.cuh in favor of libcu++'s limits #654

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.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
52 changes: 0 additions & 52 deletions common/base/include/claraparabricks/genomeworks/utils/limits.cuh
View file
Open in desktop

This file was deleted.

4 changes: 2 additions & 2 deletions cudaaligner/src/batched_device_matrices.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,7 +19,7 @@
#include "matrix_cpu.hpp"

#include <claraparabricks/genomeworks/utils/cudautils.hpp>
#include <claraparabricks/genomeworks/utils/limits.cuh>
#include <cuda/std/limits>
#include <claraparabricks/genomeworks/utils/signed_integer_utils.hpp>
#include <claraparabricks/genomeworks/utils/device_buffer.hpp>
#include <claraparabricks/genomeworks/utils/pinned_host_vector.hpp>
Expand Down Expand Up @@ -131,7 +131,7 @@ public:
{
assert(id < n_matrices_);
assert(offsets_[id + 1] - offsets_[id] >= 0);
assert(offsets_[id + 1] - offsets_[id] <= numeric_limits<int32_t>::max());
assert(offsets_[id + 1] - offsets_[id] <= cuda::std::numeric_limits<int32_t>::max());
return offsets_[id + 1] - offsets_[id];
}

Expand Down
4 changes: 2 additions & 2 deletions cudaaligner/src/hirschberg_myers_gpu.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,7 +20,7 @@
#include <claraparabricks/genomeworks/cudaaligner/aligner.hpp>
#include <claraparabricks/genomeworks/utils/cudautils.hpp>
#include <claraparabricks/genomeworks/utils/mathutils.hpp>
#include <claraparabricks/genomeworks/utils/limits.cuh>
#include <cuda/std/limits>
#include <cstring>

namespace claraparabricks
Expand Down Expand Up @@ -455,7 +455,7 @@ __device__ const char* hirschberg_myers_compute_target_mid_warp(

const int32_t target_size = (target_end - target_begin);
int32_t midpoint = 0;
nw_score_t cur_min = numeric_limits<nw_score_t>::max();
nw_score_t cur_min = cuda::std::numeric_limits<nw_score_t>::max();
for (int32_t t = threadIdx.x; t <= target_size; t += warp_size)
{
nw_score_t sum = score(t, 0) + score(target_size - t, 1);
Expand Down
5 changes: 2 additions & 3 deletions cudaaligner/src/myers_gpu.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,14 +19,13 @@

#include <claraparabricks/genomeworks/cudaaligner/aligner.hpp>
#include <claraparabricks/genomeworks/utils/signed_integer_utils.hpp>
#include <claraparabricks/genomeworks/utils/limits.cuh>
#include <claraparabricks/genomeworks/utils/mathutils.hpp>
#include <claraparabricks/genomeworks/utils/cudautils.hpp>
#include <claraparabricks/genomeworks/utils/allocator.hpp>
#include <claraparabricks/genomeworks/utils/device_buffer.hpp>

#include <cassert>
#include <climits>
#include <cuda/std/limits>
#include <vector>
#include <numeric>
#pragma GCC diagnostic push
Expand Down Expand Up @@ -429,7 +428,7 @@ __device__ int32_t myers_backtrace_banded(int8_t* path, int32_t* const path_coun
{
assert(threadIdx.x == 0);
using nw_score_t = int32_t;
GW_CONSTEXPR nw_score_t out_of_band = numeric_limits<nw_score_t>::max() - 1; // -1 to avoid integer overflow further down.
GW_CONSTEXPR nw_score_t out_of_band = cuda::std::numeric_limits<nw_score_t>::max() - 1; // -1 to avoid integer overflow further down.
assert(pv.num_rows() == score.num_rows());
assert(mv.num_rows() == score.num_rows());
assert(pv.num_cols() == score.num_cols());
Expand Down
10 changes: 5 additions & 5 deletions cudaaligner/src/ukkonen_gpu.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,7 +17,7 @@
#include "ukkonen_gpu.cuh"
#include "batched_device_matrices.cuh"
#include <claraparabricks/genomeworks/cudaaligner/cudaaligner.hpp>
#include <claraparabricks/genomeworks/utils/limits.cuh>
#include <cuda/std/limits>

#include <limits>
#include <cstdint>
Expand Down Expand Up @@ -77,7 +77,7 @@ __launch_bounds__(GW_UKKONEN_MAX_THREADS_PER_BLOCK) // Workaround for a register
if (id >= n_alignments)
return;

GW_CONSTEXPR nw_score_t max = numeric_limits<nw_score_t>::max() - 1;
GW_CONSTEXPR nw_score_t max = cuda::std::numeric_limits<nw_score_t>::max() - 1;

int32_t m = sequence_lengths_d[2 * id] + 1;
int32_t n = sequence_lengths_d[2 * id + 1] + 1;
Expand Down Expand Up @@ -153,7 +153,7 @@ __launch_bounds__(GW_UKKONEN_MAX_THREADS_PER_BLOCK) // Workaround for a register

__device__ void ukkonen_compute_score_matrix_odd(device_matrix_view<nw_score_t>& scores, int32_t kmax, int32_t k, int32_t m, int32_t n, char const* query, char const* target, int32_t max_target_query_length, int32_t p, int32_t l)
{
GW_CONSTEXPR nw_score_t max = numeric_limits<nw_score_t>::max() - 1;
GW_CONSTEXPR nw_score_t max = cuda::std::numeric_limits<nw_score_t>::max() - 1;
while (k < kmax)
{
int32_t const lmin = abs(2 * k + 1 - p);
Expand All @@ -173,7 +173,7 @@ __device__ void ukkonen_compute_score_matrix_odd(device_matrix_view<nw_score_t>&

__device__ void ukkonen_compute_score_matrix_even(device_matrix_view<nw_score_t>& scores, int32_t kmax, int32_t k, int32_t m, int32_t n, char const* query, char const* target, int32_t max_target_query_length, int32_t p, int32_t l)
{
GW_CONSTEXPR nw_score_t max = numeric_limits<nw_score_t>::max() - 1;
GW_CONSTEXPR nw_score_t max = cuda::std::numeric_limits<nw_score_t>::max() - 1;
while (k < kmax)
{
int32_t const lmin = abs(2 * k - p);
Expand All @@ -193,7 +193,7 @@ __device__ void ukkonen_compute_score_matrix_even(device_matrix_view<nw_score_t>

__device__ void ukkonen_init_score_matrix(device_matrix_view<nw_score_t>& scores, int32_t k, int32_t p)
{
GW_CONSTEXPR nw_score_t max = numeric_limits<nw_score_t>::max() - 1;
GW_CONSTEXPR nw_score_t max = cuda::std::numeric_limits<nw_score_t>::max() - 1;
while (k < scores.num_rows())
{
for (int32_t l = 0; l < scores.num_cols(); ++l)
Expand Down
6 changes: 3 additions & 3 deletions cudapoa/src/cudapoa_nw.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,9 +19,9 @@
#include "cudapoa_structs.cuh"

#include <claraparabricks/genomeworks/utils/cudautils.hpp>
#include <claraparabricks/genomeworks/utils/limits.cuh>
#include <cuda/std/limits>

#include <stdio.h>
#include <cstdio>

namespace claraparabricks
{
Expand Down Expand Up @@ -172,7 +172,7 @@ __device__ __forceinline__

static_assert(CPT == 4, "implementation currently supports only 4 cells per thread");

GW_CONSTEXPR ScoreT score_type_min_limit = numeric_limits<ScoreT>::min();
GW_CONSTEXPR ScoreT score_type_min_limit = cuda::std::numeric_limits<ScoreT>::min();

int16_t lane_idx = threadIdx.x % WARP_SIZE;
int64_t score_index;
Expand Down
6 changes: 3 additions & 3 deletions cudapoa/src/cudapoa_nw_banded.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,9 +19,9 @@
#include "cudapoa_structs.cuh"

#include <claraparabricks/genomeworks/utils/cudautils.hpp>
#include <claraparabricks/genomeworks/utils/limits.cuh>
#include <cuda/std/limits>

#include <stdio.h>
#include <cstdio>

namespace claraparabricks
{
Expand Down Expand Up @@ -199,7 +199,7 @@ __device__ __forceinline__
int32_t match_score,
int32_t rerun)
{
const ScoreT min_score_value = numeric_limits<ScoreT>::min() / 2;
constexpr ScoreT min_score_value = cuda::std::numeric_limits<ScoreT>::min() / 2;

int32_t lane_idx = threadIdx.x % WARP_SIZE;

Expand Down
6 changes: 3 additions & 3 deletions cudapoa/src/cudapoa_nw_tb_banded.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,9 +19,9 @@
#include "cudapoa_structs.cuh"

#include <claraparabricks/genomeworks/utils/cudautils.hpp>
#include <claraparabricks/genomeworks/utils/limits.cuh>
#include <cuda/std/limits>

#include <stdio.h>
#include <cstdio>

namespace claraparabricks
{
Expand Down Expand Up @@ -289,7 +289,7 @@ __device__ __forceinline__
int32_t match_score,
int32_t rerun)
{
const ScoreT min_score_value = numeric_limits<ScoreT>::min() / 2;
constexpr ScoreT min_score_value = cuda::std::numeric_limits<ScoreT>::min() / 2;

int32_t lane_idx = threadIdx.x % WARP_SIZE;

Expand Down