Индексирование CUDA не работает должным образом

Question

Я пытаюсь обработать 2D-массив с помощью PyCUDA, и мне нужны координаты x, y для каждого потока.

Этот вопрос задан и отвечает here и here, но связанные решения не работают для меня для 2D-данных, которые превышают мой размер блока. Зачем?

Вот SourceModule я использую, чтобы помочь понять это:

mod = SourceModule(""" 
    __global__ void kIndexTest(float *M, float *X, float*Y) 
    { 
    int bIdx = blockIdx.x + blockIdx.y * gridDim.x; 
    int idx = bIdx * (blockDim.x * blockDim.y) + (threadIdx.y * blockDim.x) + threadIdx.x; 

    /* this array shows me the unique thread indices */ 
    M[idx] = idx; 

    /* these arrays should capture x, y for each unique index */  
    X[idx] = (blockDim.x * blockIdx.x) + threadIdx.x; 
    Y[idx] = (blockDim.y * blockIdx.y) + threadIdx.y; 

    } 
    """) 

Я исполнения ядра, как это:

gIndexTest = mod.get_function("kIndexTest") 

dims = (8, 8) 

M = gpuarray.to_gpu(numpy.zeros(dims, dtype=numpy.float32)) 
X = gpuarray.to_gpu(numpy.zeros(dims, dtype=numpy.float32)) 
Y = gpuarray.to_gpu(numpy.zeros(dims, dtype=numpy.float32)) 

gIndexTest(M, X, Y, block=(4, 4, 1), grid=(2, 2, 1)) 

M возвращает правильный индекс для всех размеров и всех конфигурация блоков/сетки, которые я тестировал. X и Y возвращают правильные значения координат, когда размеры X и Y совпадают с размерами блока, но не возвращают то, что я ожидаю иначе. Например, вышеуказанные выходы конфигурации:

M: 
[[ 0. 1. 2. 3. 4. 5. 6. 7.] 
[ 8. 9. 10. 11. 12. 13. 14. 15.] 
[ 16. 17. 18. 19. 20. 21. 22. 23.] 
[ 24. 25. 26. 27. 28. 29. 30. 31.] 
[ 32. 33. 34. 35. 36. 37. 38. 39.] 
[ 40. 41. 42. 43. 44. 45. 46. 47.] 
[ 48. 49. 50. 51. 52. 53. 54. 55.] 
[ 56. 57. 58. 59. 60. 61. 62. 63.]] (correct) 

X: 
[[ 0. 1. 2. 3. 0. 1. 2. 3.] 
[ 0. 1. 2. 3. 0. 1. 2. 3.] 
[ 4. 5. 6. 7. 4. 5. 6. 7.] 
[ 4. 5. 6. 7. 4. 5. 6. 7.] 
[ 0. 1. 2. 3. 0. 1. 2. 3.] 
[ 0. 1. 2. 3. 0. 1. 2. 3.] 
[ 4. 5. 6. 7. 4. 5. 6. 7.] 
[ 4. 5. 6. 7. 4. 5. 6. 7.]] (not what I expect) 

Y: 
[[ 0. 0. 0. 0. 1. 1. 1. 1.] 
[ 2. 2. 2. 2. 3. 3. 3. 3.] 
[ 0. 0. 0. 0. 1. 1. 1. 1.] 
[ 2. 2. 2. 2. 3. 3. 3. 3.] 
[ 4. 4. 4. 4. 5. 5. 5. 5.] 
[ 6. 6. 6. 6. 7. 7. 7. 7.] 
[ 4. 4. 4. 4. 5. 5. 5. 5.] 
[ 6. 6. 6. 6. 7. 7. 7. 7.]] (not what I expect) 

Вот что я на самом деле ожидать от X и Y:

X: 
[[ 0. 1. 2. 3. 4. 5. 6. 7.] 
[ 0. 1. 2. 3. 4. 5. 6. 7.] 
[ 0. 1. 2. 3. 4. 5. 6. 7.] 
[ 0. 1. 2. 3. 4. 5. 6. 7.] 
[ 0. 1. 2. 3. 4. 5. 6. 7.] 
[ 0. 1. 2. 3. 4. 5. 6. 7.] 
[ 0. 1. 2. 3. 4. 5. 6. 7.] 
[ 0. 1. 2. 3. 4. 5. 6. 7.]] (only works when X dims = block dims) 

Y: 
[[ 0. 0. 0. 0. 0. 0. 0. 0.] 
[ 1. 1. 1. 1. 1. 1. 1. 1.] 
[ 2. 2. 2. 2. 2. 2. 2. 2.] 
[ 3. 3. 3. 3. 3. 3. 3. 3.] 
[ 4. 4. 4. 4. 4. 4. 4. 4.] 
[ 5. 5. 5. 5. 5. 5. 5. 5.] 
[ 6. 6. 6. 6. 6. 6. 6. 6.] 
[ 7. 7. 7. 7. 7. 7. 7. 7.]] (only works when Y dims = block dims) 

Что я не понимаю?

Вот мой deviceQuery:

Device 0: "GeForce GT 755M" 
    CUDA Driver Version/Runtime Version   7.5/6.5 
    CUDA Capability Major/Minor version number: 3.0 
    Total amount of global memory:     1024 MBytes (1073283072 bytes) 
    (2) Multiprocessors, (192) CUDA Cores/MP:  384 CUDA Cores 
    GPU Clock rate:        1085 MHz (1.09 GHz) 
    Memory Clock rate:        2500 Mhz 
    Memory Bus Width:        128-bit 
    L2 Cache Size:         262144 bytes 
    Maximum Texture Dimension Size (x,y,z)   1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096) 
    Maximum Layered 1D Texture Size, (num) layers 1D=(16384), 2048 layers 
    Maximum Layered 2D Texture Size, (num) layers 2D=(16384, 16384), 2048 layers 
    Total amount of constant memory:    65536 bytes 
    Total amount of shared memory per block:  49152 bytes 
    Total number of registers available per block: 65536 
    Warp size:          32 
    Maximum number of threads per multiprocessor: 2048 
    Maximum number of threads per block:   1024 
    Max dimension size of a thread block (x,y,z): (1024, 1024, 64) 
    Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535) 
    Maximum memory pitch:       2147483647 bytes 
    Texture alignment:        512 bytes 
    Concurrent copy and kernel execution:   Yes with 1 copy engine(s) 
    Run time limit on kernels:      Yes 
    Integrated GPU sharing Host Memory:   No 
    Support host page-locked memory mapping:  Yes 
    Alignment requirement for Surfaces:   Yes 
    Device has ECC support:      Disabled 
    Device supports Unified Addressing (UVA):  Yes 
    Device PCI Bus ID/PCI location ID:   1/0 

Answer 1

Все работает "как рекламируется". Проблема здесь в том, что вы смешиваете несовместимые схемы индексирования вместе, что приводит к непоследовательным результатам.

Если вы хотите X и Y появляться, как вы ожидали, что вам нужно будет вычислять idx по-другому:

__global__ void kIndexTest(float *M, float *X, float*Y) 
    { 
    int xidx = (blockDim.x * blockIdx.x) + threadIdx.x; 
    int yidx = (blockDim.y * blockIdx.y) + threadIdx.y; 
    int idx = (gridDim.x * blockDim.x * yidx) + xidx; 

    X[idx] = xidx; 
    Y[idx] = yidx; 
    M[idx] = idx; 
    } 

в этой схеме, xidx и yidx являются сеточными координатами х и у, и idx является глобальным индексом, все допуская основное упорядочение столбца (т. Е. X является самым быстро меняющимся размером).