← Back to branch summary

~ubuntu-branches/ubuntu/utopic/libthrust/utopic

« back to all changes in this revision

Viewing changes to system/cuda/detail/detail/launch_calculator.inl

  • Committer: Package Import Robot
  • Author(s): Andreas Beckmann
  • Date: 2013-07-10 12:57:33 UTC
  • mfrom: (1.1.4)
  • Revision ID: package-import@ubuntu.com-20130710125733-my19jic71sqsabaj
Tags: 1.7.0-1
http://bugs.debian.org/715362
* New upstream release.  (Closes: #715362)
* Update watch file.

Show diffs side-by-side

added added

removed removed

Lines of Context:
system/cuda/detail/detail/launch_calculator.inl
17
17
// do not attempt to compile this file with any other compiler
18
18
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
19
19
 
20
 
#include <thrust/system/cuda/detail/arch.h>
 
20
#include <thrust/system/cuda/detail/runtime_introspection.h>
 
21
#include <thrust/system/cuda/detail/cuda_launch_config.h>
21
22
#include <thrust/system/cuda/detail/detail/launch_closure.h>
22
23
 
23
24
namespace thrust
33
34
 
34
35
template <typename Closure>
35
36
launch_calculator<Closure>::launch_calculator(void)
36
 
  : properties(arch::device_properties()),
 
37
  : properties(device_properties()),
37
38
    attributes(closure_attributes<Closure>())
38
39
{}
39
40
  
40
41
template <typename Closure>
41
 
launch_calculator<Closure>::launch_calculator(const arch::device_properties_t& properties, const arch::function_attributes_t& attributes)
 
42
launch_calculator<Closure>::launch_calculator(const device_properties_t& properties, const function_attributes_t& attributes)
42
43
  : properties(properties),
43
44
    attributes(attributes)
44
45
{}
45
46
 
46
47
template <typename Closure>
47
 
thrust::tuple<size_t,size_t,size_t> launch_calculator<Closure>::with_variable_block_size(void)
48
 
{
49
 
  thrust::pair<size_t, size_t> config = arch::default_block_configuration(properties, attributes);
 
48
  template <typename UnaryFunction>
 
49
thrust::pair<size_t, size_t> launch_calculator<Closure>::default_block_configuration(UnaryFunction block_size_to_smem_size) const
 
50
{
 
51
  // choose a block size
 
52
  std::size_t num_threads_per_block = block_size_with_maximum_potential_occupancy(attributes, properties, block_size_to_smem_size);
 
53
 
 
54
  // choose a subscription rate
 
55
  std::size_t num_blocks_per_multiprocessor = properties.maxThreadsPerMultiProcessor / num_threads_per_block;
 
56
 
 
57
  return thrust::make_pair(num_threads_per_block, num_blocks_per_multiprocessor);
 
58
}
 
59
 
 
60
 
 
61
template <typename Closure>
 
62
thrust::pair<size_t, size_t> launch_calculator<Closure>::default_block_configuration(void) const
 
63
{
 
64
  // choose a block size
 
65
  std::size_t num_threads_per_block = block_size_with_maximum_potential_occupancy(attributes, properties);
 
66
 
 
67
  // choose a subscription rate
 
68
  std::size_t num_blocks_per_multiprocessor = properties.maxThreadsPerMultiProcessor / num_threads_per_block;
 
69
 
 
70
  return thrust::make_pair(num_threads_per_block, num_blocks_per_multiprocessor);
 
71
}
 
72
 
 
73
template <typename Closure>
 
74
thrust::tuple<size_t,size_t,size_t> launch_calculator<Closure>::with_variable_block_size(void) const
 
75
{
 
76
  thrust::pair<size_t, size_t> config = default_block_configuration();
50
77
  return thrust::tuple<size_t,size_t,size_t>(config.second * properties.multiProcessorCount, config.first, 0);
51
78
}
52
79
 
53
80
template <typename Closure>
54
81
  template <typename UnaryFunction>
55
 
thrust::tuple<size_t,size_t,size_t> launch_calculator<Closure>::with_variable_block_size(UnaryFunction block_size_to_smem_size)
 
82
thrust::tuple<size_t,size_t,size_t> launch_calculator<Closure>::with_variable_block_size(UnaryFunction block_size_to_smem_size) const
56
83
{
57
 
  thrust::pair<size_t, size_t> config = arch::default_block_configuration(properties, attributes, block_size_to_smem_size);
 
84
  thrust::pair<size_t, size_t> config = default_block_configuration(block_size_to_smem_size);
58
85
  return thrust::tuple<size_t,size_t,size_t>(config.second * properties.multiProcessorCount, config.first, block_size_to_smem_size(config.first));
59
86
}
60
87
  
61
88
template <typename Closure>
62
 
thrust::tuple<size_t,size_t,size_t> launch_calculator<Closure>::with_variable_block_size_available_smem(void)
 
89
thrust::tuple<size_t,size_t,size_t> launch_calculator<Closure>::with_variable_block_size_available_smem(void) const
63
90
{
64
 
  thrust::pair<size_t, size_t> config = arch::default_block_configuration(properties, attributes);
65
 
  size_t smem_per_block = arch::proportional_smem_allocation(properties, attributes, config.second);
 
91
  thrust::pair<size_t, size_t> config = default_block_configuration();
 
92
  size_t smem_per_block = proportional_smem_allocation(properties, attributes, config.second);
66
93
  return thrust::tuple<size_t,size_t,size_t>(config.second * properties.multiProcessorCount, config.first, smem_per_block);
67
94
}
68
95