在我的代码中，我在全局内存中有一个静态分配的数组（即，使用分配了__device__），而我想使用进行排序thrust::sort，这是行不通的。关于该主题的所有示例都使用CUDA运行时分配的数组（使用cudaMalloc）。有什么办法可以对静态分配的数组进行排序？
我想这与无法从主机访问的静态分配内存有关。使用cudaMalloc-allocated数组，可以正常工作。但是，我想避免使用这种类型的分配，因为静态分配可以更轻松地从设备代码访问数据（不是吗？）。
#include <stdio.h>
#include <thrust/device_ptr.h>
#include <thrust/sort.h>
#define N 4
typedef struct element {
int key;
int value;
__host__ __device__ bool operator<(element e) const
{ return key > e.key; }
} element;
__device__ element array[N];
__global__ void init() {
for (int i = 0; i < N; ++i) {
array[N - i - 1].key = i;
}
}
__global__ void print_array() {
for (int i = 0; i < N; ++i) {
printf("%d ", array[i].key);
}
printf("\n");
}
int main(void) {
thrust::device_ptr<element> array_first(array);
init<<<1,1>>>();
printf("unsorted: ");
print_array<<<1, 1>>>();
cudaDeviceSynchronize();
thrust::sort(array_first, array_first + N);
printf("sorted: ");
print_array<<<1, 1>>>();
cudaDeviceSynchronize();
}
用于从函数cudaGetSymbolAddress获取array变量的地址__host__：
void* array_ptr = 0;
cudaGetSymbolAddress(&array_ptr, array);
thrust::device_ptr<element> array_first(reinterpret_cast<element*>(array_ptr));
这是完整的程序：
#include <stdio.h>
#include <thrust/device_ptr.h>
#include <thrust/sort.h>
#define N 4
typedef struct element {
int key;
int value;
__host__ __device__ bool operator<(element e) const
{ return key > e.key; }
} element;
__device__ element array[N];
__global__ void init() {
for (int i = 0; i < N; ++i) {
array[N - i - 1].key = i;
}
}
__global__ void print_array() {
for (int i = 0; i < N; ++i) {
printf("%d ", array[i].key);
}
printf("\n");
}
int main(void) {
cudaError_t error;
void* array_ptr = 0;
if(error = cudaGetSymbolAddress(&array_ptr, array))
{
throw thrust::system_error(error, thrust::cuda_category());
}
thrust::device_ptr<element> array_first(reinterpret_cast<element*>(array_ptr));
init<<<1,1>>>();
printf("unsorted: ");
print_array<<<1, 1>>>();
if(error = cudaDeviceSynchronize())
{
throw thrust::system_error(error, thrust::cuda_category());
}
thrust::sort(array_first, array_first + N);
if(error = cudaDeviceSynchronize())
{
throw thrust::system_error(error, thrust::cuda_category());
}
printf("sorted: ");
print_array<<<1, 1>>>();
if(error = cudaDeviceSynchronize())
{
throw thrust::system_error(error, thrust::cuda_category());
}
return 0;
}
这是我系统上的输出：
$ nvcc test.cu -run
unsorted: 3 2 1 0
sorted: 3 2 1 0
排序后的输出与未排序后的输出相同，但是考虑到数据的生成方式和的定义，我认为这是有意的element::operator<。
本文收集自互联网，转载请注明来源。
如有侵权，请联系[email protected] 删除。
我来说两句