mirror of
https://gitlab.kitware.com/vtk/vtk-m
synced 2024-09-19 18:45:43 +00:00
Merge topic 'reenable_cuda_texture_fastpath'
b2cd41d7 Fix ArrayPortalFromThrust to re-enable texture memory fast path. Acked-by: Kitware Robot <kwrobot@kitware.com> Merge-request: !330
This commit is contained in:
commit
6ad3f115ab
@ -95,6 +95,8 @@ template<> struct UseMultipleScalarTextureLoads< const vtkm::Vec<vtkm::Float64,4
|
||||
template<typename T, typename Enable = void>
|
||||
struct load_through_texture
|
||||
{
|
||||
static const vtkm::IdComponent WillUseTexture = 0;
|
||||
|
||||
__device__
|
||||
static T get(const thrust::system::cuda::pointer<const T>& data)
|
||||
{
|
||||
@ -106,8 +108,11 @@ struct load_through_texture
|
||||
|
||||
// this T type is valid to be loaded through a single texture memory fetch
|
||||
template<typename T>
|
||||
struct load_through_texture<T, typename ::boost::enable_if< typename UseScalarTextureLoad<T>::type >::type >
|
||||
struct load_through_texture<T, typename ::boost::enable_if< typename UseScalarTextureLoad<const T>::type >::type >
|
||||
{
|
||||
|
||||
static const vtkm::IdComponent WillUseTexture = 1;
|
||||
|
||||
__device__
|
||||
static T get(const thrust::system::cuda::pointer<const T>& data)
|
||||
{
|
||||
@ -122,8 +127,10 @@ struct load_through_texture<T, typename ::boost::enable_if< typename UseScalarTe
|
||||
|
||||
// this T type is valid to be loaded through a single vec texture memory fetch
|
||||
template<typename T>
|
||||
struct load_through_texture<T, typename ::boost::enable_if< typename UseVecTextureLoads<T>::type >::type >
|
||||
struct load_through_texture<T, typename ::boost::enable_if< typename UseVecTextureLoads<const T>::type >::type >
|
||||
{
|
||||
static const vtkm::IdComponent WillUseTexture = 1;
|
||||
|
||||
__device__
|
||||
static T get(const thrust::system::cuda::pointer<const T>& data)
|
||||
{
|
||||
@ -188,8 +195,10 @@ struct load_through_texture<T, typename ::boost::enable_if< typename UseVecTextu
|
||||
|
||||
//this T type is valid to be loaded through multiple texture memory fetches
|
||||
template<typename T>
|
||||
struct load_through_texture<T, typename ::boost::enable_if< typename UseMultipleScalarTextureLoads<T>::type >::type >
|
||||
struct load_through_texture<T, typename ::boost::enable_if< typename UseMultipleScalarTextureLoads<const T>::type >::type >
|
||||
{
|
||||
static const vtkm::IdComponent WillUseTexture = 1;
|
||||
|
||||
typedef typename boost::remove_const<T>::type NonConstT;
|
||||
|
||||
__device__
|
||||
|
@ -32,3 +32,8 @@ if (VTKm_ENABLE_CUDA)
|
||||
endif()
|
||||
|
||||
vtkm_declare_headers(CUDA ${headers} TESTABLE ${VTKm_ENABLE_CUDA})
|
||||
|
||||
#-----------------------------------------------------------------------------
|
||||
if (VTKm_ENABLE_CUDA)
|
||||
add_subdirectory(testing)
|
||||
endif()
|
||||
|
26
vtkm/exec/cuda/internal/testing/CMakeLists.txt
Normal file
26
vtkm/exec/cuda/internal/testing/CMakeLists.txt
Normal file
@ -0,0 +1,26 @@
|
||||
##=============================================================================
|
||||
##
|
||||
## Copyright (c) Kitware, Inc.
|
||||
## All rights reserved.
|
||||
## See LICENSE.txt for details.
|
||||
##
|
||||
## This software is distributed WITHOUT ANY WARRANTY; without even
|
||||
## the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
|
||||
## PURPOSE. See the above copyright notice for more information.
|
||||
##
|
||||
## Copyright 2014 Sandia Corporation.
|
||||
## Copyright 2014 UT-Battelle, LLC.
|
||||
## Copyright 2014 Los Alamos National Security.
|
||||
##
|
||||
## Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
|
||||
## the U.S. Government retains certain rights in this software.
|
||||
## Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
|
||||
## Laboratory (LANL), the U.S. Government retains certain rights in
|
||||
## this software.
|
||||
##
|
||||
##=============================================================================
|
||||
|
||||
set(unit_tests
|
||||
UnitTestTextureMemorySupport.cu
|
||||
)
|
||||
vtkm_unit_tests(CUDA SOURCES ${unit_tests})
|
@ -0,0 +1,88 @@
|
||||
//============================================================================
|
||||
// Copyright (c) Kitware, Inc.
|
||||
// All rights reserved.
|
||||
// See LICENSE.txt for details.
|
||||
// This software is distributed WITHOUT ANY WARRANTY; without even
|
||||
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
|
||||
// PURPOSE. See the above copyright notice for more information.
|
||||
//
|
||||
// Copyright 2015 Sandia Corporation.
|
||||
// Copyright 2015 UT-Battelle, LLC.
|
||||
// Copyright 2015 Los Alamos National Security.
|
||||
//
|
||||
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
|
||||
// the U.S. Government retains certain rights in this software.
|
||||
//
|
||||
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
|
||||
// Laboratory (LANL), the U.S. Government retains certain rights in
|
||||
// this software.
|
||||
//============================================================================
|
||||
|
||||
#include <vtkm/exec/cuda/internal/ArrayPortalFromThrust.h>
|
||||
#include <vtkm/cont/testing/Testing.h>
|
||||
|
||||
namespace {
|
||||
|
||||
struct customType { };
|
||||
|
||||
void TestScalarTextureLoad()
|
||||
{
|
||||
using namespace vtkm::exec::cuda::internal;
|
||||
typedef load_through_texture< vtkm::Float32 > f;
|
||||
typedef load_through_texture< vtkm::Int32 > i;
|
||||
typedef load_through_texture< vtkm::UInt8 > ui;
|
||||
|
||||
typedef load_through_texture< customType > ct;
|
||||
|
||||
VTKM_TEST_ASSERT( f::WillUseTexture == 1, "Float32 can be loaded through texture memory" );
|
||||
VTKM_TEST_ASSERT( i::WillUseTexture == 1, "Int32 can be loaded through texture memory" );
|
||||
VTKM_TEST_ASSERT( ui::WillUseTexture == 1, "Unsigned Int8 can be loaded through texture memory" );
|
||||
VTKM_TEST_ASSERT( ct::WillUseTexture == 0, "Custom Types can't be loaded through texture memory" );
|
||||
|
||||
}
|
||||
|
||||
void TestVecTextureLoad()
|
||||
{
|
||||
using namespace
|
||||
vtkm::exec::cuda::internal;
|
||||
typedef load_through_texture< vtkm::Vec<vtkm::UInt32,3> > ui32_3;
|
||||
typedef load_through_texture< vtkm::Vec<vtkm::Float32,3> > f32_3;
|
||||
typedef load_through_texture< vtkm::Vec<vtkm::UInt8,3> > ui8_3;
|
||||
typedef load_through_texture< vtkm::Vec<vtkm::Float64,3> > f64_3;
|
||||
|
||||
typedef load_through_texture< vtkm::Vec<vtkm::UInt32,4> > ui32_4;
|
||||
typedef load_through_texture< vtkm::Vec<vtkm::Float32,4> > f32_4;
|
||||
typedef load_through_texture< vtkm::Vec<vtkm::UInt8,4> > ui8_4;
|
||||
typedef load_through_texture< vtkm::Vec<vtkm::Float64,4> > f64_4;
|
||||
|
||||
typedef load_through_texture< vtkm::Vec<customType, 3> > ct_3;
|
||||
typedef load_through_texture< vtkm::Vec<customType, 4> > ct_4;
|
||||
|
||||
|
||||
VTKM_TEST_ASSERT( ui32_3::WillUseTexture == 1, "Can be loaded through texture loads");
|
||||
VTKM_TEST_ASSERT( f32_3::WillUseTexture == 1, "Can be loaded through texture loads");
|
||||
VTKM_TEST_ASSERT( ui8_3::WillUseTexture == 1, "Can be loaded through texture loads");
|
||||
VTKM_TEST_ASSERT( f64_3::WillUseTexture == 1, "Can be loaded through texture loads");
|
||||
|
||||
VTKM_TEST_ASSERT( ui32_4::WillUseTexture == 1, "Can be loaded through texture loads");
|
||||
VTKM_TEST_ASSERT( f32_4::WillUseTexture == 1, "Can be loaded through texture loads");
|
||||
VTKM_TEST_ASSERT( ui8_4::WillUseTexture == 1, "Can be loaded through texture loads");
|
||||
VTKM_TEST_ASSERT( f64_4::WillUseTexture == 1, "Can be loaded through texture loads");
|
||||
|
||||
VTKM_TEST_ASSERT( ct_4::WillUseTexture == 0, "Can't be loaded through texture loads");
|
||||
VTKM_TEST_ASSERT( ct_4::WillUseTexture == 0, "Can't be loaded through texture loads");
|
||||
}
|
||||
|
||||
|
||||
} // namespace
|
||||
|
||||
void TestTextureMemorySupport()
|
||||
{
|
||||
TestScalarTextureLoad();
|
||||
TestVecTextureLoad();
|
||||
}
|
||||
|
||||
int UnitTestTextureMemorySupport(int, char *[])
|
||||
{
|
||||
return vtkm::cont::testing::Testing::Run( TestTextureMemorySupport );
|
||||
}
|
Loading…
Reference in New Issue
Block a user