Commit b7e63718 authored by Robert Maynard's avatar Robert Maynard
Browse files

Correct issues found be enabling more CUDA warnings.

parent 4ca2522c
......@@ -2267,12 +2267,12 @@ static inline VTKM_EXEC_CONT vtkm::Float32 RemainderQuotient(vtkm::Float32 numer
{
int iQuotient;
#ifdef VTKM_CUDA
const vtkm::Float64 result =
const vtkm::Float32 result =
VTKM_CUDA_MATH_FUNCTION_32(remquo)(numerator, denominator, &iQuotient);
#else
const vtkm::Float32 result = std::remquo(numerator, denominator, &iQuotient);
#endif
quotient = iQuotient;
quotient = static_cast<QType>(iQuotient);
return result;
}
template <typename QType>
......@@ -2287,7 +2287,7 @@ static inline VTKM_EXEC_CONT vtkm::Float64 RemainderQuotient(vtkm::Float64 numer
#else
const vtkm::Float64 result = std::remquo(numerator, denominator, &iQuotient);
#endif
quotient = iQuotient;
quotient = static_cast<QType>(iQuotient);
return result;
}
......
......@@ -922,12 +922,12 @@ static inline VTKM_EXEC_CONT vtkm::Float32 RemainderQuotient(vtkm::Float32 numer
{
int iQuotient;
#ifdef VTKM_CUDA
const vtkm::Float64 result =
const vtkm::Float32 result =
VTKM_CUDA_MATH_FUNCTION_32(remquo)(numerator, denominator, &iQuotient);
#else
const vtkm::Float32 result = std::remquo(numerator, denominator, &iQuotient);
#endif
quotient = iQuotient;
quotient = static_cast<QType>(iQuotient);
return result;
}
template <typename QType>
......@@ -942,7 +942,7 @@ static inline VTKM_EXEC_CONT vtkm::Float64 RemainderQuotient(vtkm::Float64 numer
#else
const vtkm::Float64 result = std::remquo(numerator, denominator, &iQuotient);
#endif
quotient = iQuotient;
quotient = static_cast<QType>(iQuotient);
return result;
}
......
......@@ -77,6 +77,11 @@ namespace cuda
namespace internal
{
#if (defined(VTKM_GCC) || defined(VTKM_CLANG))
#pragma GCC diagnostic push
#pragma GCC diagnostic ignored "-Wunused-parameter"
#endif
static __global__ void DetermineProperXGridSize(vtkm::UInt32 desired_size,
vtkm::UInt32* actual_size)
{
......@@ -171,6 +176,10 @@ __global__ void SumExclusiveScan(T a, T b, T result, BinaryOperationType binary_
result = binary_op(a, b);
}
#if (defined(VTKM_GCC) || defined(VTKM_CLANG))
#pragma GCC diagnostic pop
#endif
inline void compute_block_size(dim3 rangeMax, dim3 blockSize3d, dim3& gridSize3d)
{
gridSize3d.x = (rangeMax.x % blockSize3d.x != 0) ? (rangeMax.x / blockSize3d.x + 1)
......
......@@ -180,22 +180,24 @@ void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::CopyToControl(const voi
cudaStreamPerThread));
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::UsingForRead(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::UsingForRead(
const void* vtkmNotUsed(controlPtr),
const void* executionPtr,
vtkm::UInt64 numBytes) const
{
CudaAllocator::PrepareForInput(executionPtr, static_cast<size_t>(numBytes));
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::UsingForWrite(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::UsingForWrite(
const void* vtkmNotUsed(controlPtr),
const void* executionPtr,
vtkm::UInt64 numBytes) const
{
CudaAllocator::PrepareForOutput(executionPtr, static_cast<size_t>(numBytes));
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::UsingForReadWrite(
const void* controlPtr,
const void* vtkmNotUsed(controlPtr),
const void* executionPtr,
vtkm::UInt64 numBytes) const
{
......
......@@ -44,7 +44,7 @@ inline vtkm::exec::cuda::internal::IteratorFromArrayPortal<PortalType> IteratorE
const PortalType& portal)
{
vtkm::exec::cuda::internal::IteratorFromArrayPortal<PortalType> iterator(portal);
iterator += static_cast<std::size_t>(portal.GetNumberOfValues());
iterator += static_cast<std::ptrdiff_t>(portal.GetNumberOfValues());
return iterator;
}
......
......@@ -35,6 +35,11 @@ namespace internal
namespace detail
{
#if (defined(VTKM_GCC) || defined(VTKM_CLANG))
#pragma GCC diagnostic push
#pragma GCC diagnostic ignored "-Wunused-parameter"
#endif
template <typename VirtualDerivedType>
__global__ void ConstructVirtualObjectKernel(VirtualDerivedType* deviceObject,
const VirtualDerivedType* targetObject)
......@@ -56,6 +61,10 @@ __global__ void DeleteVirtualObjectKernel(VirtualDerivedType* deviceObject)
deviceObject->~VirtualDerivedType();
}
#if (defined(VTKM_GCC) || defined(VTKM_CLANG))
#pragma GCC diagnostic pop
#endif
} // detail
template <typename VirtualDerivedType>
......
Supports Markdown
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