Shuenhoy/vtk-m
1
0
Fork 0
mirror of https://gitlab.kitware.com/vtk/vtk-m synced 2024-09-16 17:22:55 +00:00
Code Issues 2 Actions Packages Projects Releases Wiki Activity

Fix ArrayPortalFromThrust to re-enable texture memory fast path.

Browse Source
This commit is contained in:
Robert Maynard 2016-01-25 12:53:33 -05:00
parent bc7f996610
commit b2cd41d765
4 changed files with 131 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

15
vtkm/exec/cuda/internal/ArrayPortalFromThrust.h
View File

@ -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__

5
vtkm/exec/cuda/internal/CMakeLists.txt
View File

@ -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
View 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})

88
vtkm/exec/cuda/internal/testing/UnitTestTextureMemorySupport.cu Normal file
View File

@ -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 );
}