Рецепт для копирования данных 1D с помощью cudaMemcpy2D - PullRequest
Новая
       14

Рецепт для копирования данных 1D с помощью cudaMemcpy2D

0 голосов
/ 14 июня 2019

Если у одного есть два непрерывных диапазона памяти устройства, можно скопировать память из одного в другое, используя cudaMemcpy. 

   double* source = ...
   double* dest = ...
   cudaMemcpy(dest, source, N, cudaMemcpyDeviceToDevice);

Теперь предположим, что я хочу скопировать источник в dest, но каждые 2 или 3 элемента соответственно. Это dest[0] = source[0], dest[3] = source[2], dest[6] = source[4], .... Конечно, одна равнина cudaMemcpy не может этого сделать.

Интуитивно, cudaMemcpy2D должен быть в состоянии выполнить работу, потому что "выделенные элементы можно увидеть как столбец в большем массиве". Но cudaMemcpy2D имеет много входных параметров, которые трудно интерпретировать в этом контексте, например pitch.

Например, я могу использовать cudaMemcpy2D для воспроизведения случая, когда оба шага равны 1. 

    cudaMemcpy2D(dest, 1, source, 1, 1, n*sizeof(T), cudaMemcpyDeviceToHost);

Но я не могу понять общий случай, с dest_stride и source_stride отличием от 1.

Есть ли способ скопировать пошаговые данные в пошаговые данные с помощью cudaMemcpy2D? В каком порядке я должен поместить известную информацию о макете ?, а именно, с точки зрения двух шагов и sizeof(T). 

    cudaMemcpy2D(dest, ??, source, ???, ????, ????, cudaMemcpyDeviceToHost);

Ответы [ 2 ]

2 голосов
/ 14 июня 2019

Да, это можно сделать.Это легче проиллюстрировать в коде, чем словами: 

#include <iostream>

int main()
{
    const size_t swidth = 2;
    const size_t sheight = 4;
    size_t spitch = swidth * sizeof(int);
    int source[swidth * sheight] = { 0, 1, 2, 3, 4, 5, 6, 7 }; 


    const size_t dwidth = 3;
    const size_t dheight = 4;
    size_t dpitch = dwidth * sizeof(int);
    int dest[dwidth * dheight] = { -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1 };

    const size_t cwidth = 1 * sizeof(int);
    const size_t cheight = 3;

    int* source_d; cudaMalloc(&source_d, spitch * sheight);
    cudaMemcpy(source_d, &source[0], spitch * sheight, cudaMemcpyHostToDevice);
    cudaMemcpy2D(&dest[0], dpitch, source_d, spitch, cwidth, cheight, cudaMemcpyDeviceToHost);

    for(int i=0; i < 12; i++) std::cout << i << " " << dest[i] << std::endl;

    return 0;
}

, который делает это: 

$ nvcc -std=c++11 -arch=sm_52 -o strided_copy strided_copy.cu 
$ cuda-memcheck ./strided_copy
========= CUDA-MEMCHECK
0 0
1 -1
2 -1
3 2
4 -1
5 -1
6 4
7 -1
8 -1
9 -1
10 -1
11 -1
========= ERROR SUMMARY: 0 errors

По сути, вы копируете ширину 4 байта (целое число) сшаг 8 байтов (два целых) в пункт назначения с шагом 12 байтов (три целых).Я скопировал только три rwos, чтобы было понятно, как работает аргумент строки.Отрегулируйте размер элемента копирования и шагов и т. Д. По вкусу.

1 голос
/ 14 июня 2019

Универсальная функция для такой расширенной копии может выглядеть примерно так: 

void cudaMemcpyStrided(
        void *dst, int dstStride, 
        void *src, int srcStride, 
        int numElements, int elementSize, int kind) {
    int srcPitchInBytes = srcStride * elementSize;
    int dstPitchInBytes = dstStride * elementSize;
    int width = 1 * elementSize;
    int height = numElements;
    cudaMemcpy2D(
        dst, dstPitchInBytes, 
        src, srcPitchInBytes, 
        width, height,
        kind);
}

А для вашего примера ее можно назвать 

cudaMemcpyStrided(dest, 3, source, 2, 3, sizeof(double), cudaMemcpyDeviceToDevice);

"Примерно", потому что я только что перевел его на лету из кода (на основе Java / JCuda), с которым я его тестировал: 

import static jcuda.runtime.JCuda.cudaMemcpy2D;

import java.util.Arrays;
import java.util.Locale;

import jcuda.Pointer;
import jcuda.Sizeof;
import jcuda.runtime.cudaMemcpyKind;

public class JCudaStridedMemcopy {
    public static void main(String[] args) {

        int dstLength = 9;
        int srcLength = 6;
        int dstStride = 3;
        int srcStride = 2;
        int numElements = 3;
        runExample(dstLength, dstStride, srcLength, srcStride, numElements);

        dstLength = 9;
        srcLength = 12;
        dstStride = 3;
        srcStride = 4;
        numElements = 3;
        runExample(dstLength, dstStride, srcLength, srcStride, numElements);

        dstLength = 18;
        srcLength = 12;
        dstStride = 3;
        srcStride = 2;
        numElements = 6;
        runExample(dstLength, dstStride, srcLength, srcStride, numElements);

    }

    private static void runExample(int dstLength, int dstStride, int srcLength, int srcStride, int numElements) {
        double dst[] = new double[dstLength];
        double src[] = new double[srcLength];
        for (int i = 0; i < src.length; i++) {
            src[i] = i;
        }

        cudaMemcpyStrided(dst, dstStride, src, srcStride, numElements);

        System.out.println("Copy " + numElements + " elements");
        System.out.println("  to   array with length " + dstLength + ", with a stride of " + dstStride);
        System.out.println("  from array with length " + srcLength + ", with a stride of " + srcStride);

        System.out.println("");

        System.out.println("Destination:");
        System.out.println(toString2D(dst, dstStride));
        System.out.println("Flat: " + Arrays.toString(dst));

        System.out.println("");

        System.out.println("Source:");
        System.out.println(toString2D(src, srcStride));
        System.out.println("Flat: " + Arrays.toString(src));

        System.out.println("");
        System.out.println("Done");
        System.out.println("");

    }

    private static void cudaMemcpyStrided(double dst[], int dstStride, double src[], int srcStride, int numElements) {
        long srcPitchInBytes = srcStride * Sizeof.DOUBLE;
        long dstPitchInBytes = dstStride * Sizeof.DOUBLE;
        long width = 1 * Sizeof.DOUBLE;
        long height = numElements;
        cudaMemcpy2D(
                Pointer.to(dst), dstPitchInBytes, 
                Pointer.to(src), srcPitchInBytes, 
                width, height,
                cudaMemcpyKind.cudaMemcpyHostToHost);
    }

    public static String toString2D(double[] a, long columns) {
        String format = "%4.1f ";
        ;
        StringBuilder sb = new StringBuilder();
        for (int i = 0; i < a.length; i++) {
            if (i > 0 && i % columns == 0) {
                sb.append("\n");
            }
            sb.append(String.format(Locale.ENGLISH, format, a[i]));
        }
        return sb.toString();
    }
}

Чтобы дать представление о том, что делает функция, на основе примеров / тестаслучаи, вот вывод: 

Copy 3 elements
  to   array with length 9, with a stride of 3
  from array with length 6, with a stride of 2

Destination:
 0.0  0.0  0.0 
 2.0  0.0  0.0 
 4.0  0.0  0.0 
Flat: [0.0, 0.0, 0.0, 2.0, 0.0, 0.0, 4.0, 0.0, 0.0]

Source:
 0.0  1.0 
 2.0  3.0 
 4.0  5.0 
Flat: [0.0, 1.0, 2.0, 3.0, 4.0, 5.0]

Done

Copy 3 elements
  to   array with length 9, with a stride of 3
  from array with length 12, with a stride of 4

Destination:
 0.0  0.0  0.0 
 4.0  0.0  0.0 
 8.0  0.0  0.0 
Flat: [0.0, 0.0, 0.0, 4.0, 0.0, 0.0, 8.0, 0.0, 0.0]

Source:
 0.0  1.0  2.0  3.0 
 4.0  5.0  6.0  7.0 
 8.0  9.0 10.0 11.0 
Flat: [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0]

Done

Copy 6 elements
  to   array with length 18, with a stride of 3
  from array with length 12, with a stride of 2

Destination:
 0.0  0.0  0.0 
 2.0  0.0  0.0 
 4.0  0.0  0.0 
 6.0  0.0  0.0 
 8.0  0.0  0.0 
10.0  0.0  0.0 
Flat: [0.0, 0.0, 0.0, 2.0, 0.0, 0.0, 4.0, 0.0, 0.0, 6.0, 0.0, 0.0, 8.0, 0.0, 0.0, 10.0, 0.0, 0.0]

Source:
 0.0  1.0 
 2.0  3.0 
 4.0  5.0 
 6.0  7.0 
 8.0  9.0 
10.0 11.0 
Flat: [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0]

Done
...