Вопросы Теги

Первый параметр в инструкции shuffle war в описании CUDA 9 [дубликат]

Я думаю, что проблема - это либо неправильное имя макета, либо недопустимый путь к файлу макета.

для IntelliJ, вы можете создать каталог ресурсов и разместить там файлы макета. 

FXMLLoader loader = new FXMLLoader();
loader.setLocation(getClass().getResource("/sample.fxml"));
rootLayout = loader.load();
1
задан ztdep 1 June 2018 в 08:30
поделиться

1 ответ

В дополнение к простому имени, версии _sync для функций shuffle warp также имеют другой прототип, как указано в документации . Первый параметр - это параметр маски.

Кажется, вы пытаетесь использовать обе функции одинаково: 

     temp+=__shfl_down(temp, offset,32);

temp+=__shfl_down_sync(temp, offset,32);

, но это неверно. Чтобы использовать версию _sync аналогичным образом, вы должны сделать: 

temp+=__shfl_down_sync(0xFFFFFFFF, temp, offset,32);

Когда я делаю это изменение, ваш код работает правильно для меня: 

#include <stdio.h>

__global__ void shufledown1(double* a, double *b,double *c, int N)
{
    double  temp = 2.0;
    __syncthreads();

   for (int offset = 32/2; offset > 0; offset /= 2){
         temp+=__shfl_down(temp, offset,32);
   }
    printf("%d %f %d \n",threadIdx.x ,temp,blockDim.x * gridDim.x);
}
__global__ void shufledown2(double* a, double *b,double *c, int N)
{
    double  temp = 2.0;
    __syncthreads();

   for (int offset = 32/2; offset > 0; offset /= 2){
       temp+=__shfl_down_sync(0xFFFFFFFF, temp, offset,32);
   }
    printf("%d %f %d \n",threadIdx.x ,temp,blockDim.x * gridDim.x);
}


int main(){
    double *a = NULL, *b = NULL, *c = NULL;
    shufledown1<<<1,64>>>(a, b, c, 0);
    cudaDeviceSynchronize();
    shufledown2<<<1,64>>>(a, b, c, 0);
    cudaDeviceSynchronize();
}
$ nvcc -arch=sm_60 -o t1358 t1358.cu
t1358.cu(9): warning: function "__shfl_down(double, unsigned int, int)"
/usr/local/cuda/bin/..//include/sm_30_intrinsics.hpp(453): here was declared deprecated ("__shfl_down() is deprecated in favor of __shfl_down_sync() and may be removed in a future release (Use -Wno-deprecated-declarations to suppress this warning).")

ptxas /tmp/tmpxft_000045b6_00000000-5_t1358.ptx, line 49; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/tmpxft_000045b6_00000000-5_t1358.ptx, line 52; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/tmpxft_000045b6_00000000-5_t1358.ptx, line 63; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/tmpxft_000045b6_00000000-5_t1358.ptx, line 66; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/tmpxft_000045b6_00000000-5_t1358.ptx, line 77; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/tmpxft_000045b6_00000000-5_t1358.ptx, line 80; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/tmpxft_000045b6_00000000-5_t1358.ptx, line 91; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/tmpxft_000045b6_00000000-5_t1358.ptx, line 94; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/tmpxft_000045b6_00000000-5_t1358.ptx, line 105; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/tmpxft_000045b6_00000000-5_t1358.ptx, line 108; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
$ ./t1358
0 64.000000 64
1 64.000000 64
2 64.000000 64
3 64.000000 64
4 64.000000 64
5 64.000000 64
6 64.000000 64
7 64.000000 64
8 64.000000 64
9 64.000000 64
10 64.000000 64
11 64.000000 64
12 64.000000 64
13 64.000000 64
14 64.000000 64
15 64.000000 64
16 64.000000 64
17 64.000000 64
18 64.000000 64
19 64.000000 64
20 64.000000 64
21 64.000000 64
22 64.000000 64
23 64.000000 64
24 64.000000 64
25 64.000000 64
26 64.000000 64
27 64.000000 64
28 64.000000 64
29 64.000000 64
30 64.000000 64
31 64.000000 64
32 64.000000 64
33 64.000000 64
34 64.000000 64
35 64.000000 64
36 64.000000 64
37 64.000000 64
38 64.000000 64
39 64.000000 64
40 64.000000 64
41 64.000000 64
42 64.000000 64
43 64.000000 64
44 64.000000 64
45 64.000000 64
46 64.000000 64
47 64.000000 64
48 64.000000 64
49 64.000000 64
50 64.000000 64
51 64.000000 64
52 64.000000 64
53 64.000000 64
54 64.000000 64
55 64.000000 64
56 64.000000 64
57 64.000000 64
58 64.000000 64
59 64.000000 64
60 64.000000 64
61 64.000000 64
62 64.000000 64
63 64.000000 64
0 64.000000 64
1 64.000000 64
2 64.000000 64
3 64.000000 64
4 64.000000 64
5 64.000000 64
6 64.000000 64
7 64.000000 64
8 64.000000 64
9 64.000000 64
10 64.000000 64
11 64.000000 64
12 64.000000 64
13 64.000000 64
14 64.000000 64
15 64.000000 64
16 64.000000 64
17 64.000000 64
18 64.000000 64
19 64.000000 64
20 64.000000 64
21 64.000000 64
22 64.000000 64
23 64.000000 64
24 64.000000 64
25 64.000000 64
26 64.000000 64
27 64.000000 64
28 64.000000 64
29 64.000000 64
30 64.000000 64
31 64.000000 64
32 64.000000 64
33 64.000000 64
34 64.000000 64
35 64.000000 64
36 64.000000 64
37 64.000000 64
38 64.000000 64
39 64.000000 64
40 64.000000 64
41 64.000000 64
42 64.000000 64
43 64.000000 64
44 64.000000 64
45 64.000000 64
46 64.000000 64
47 64.000000 64
48 64.000000 64
49 64.000000 64
50 64.000000 64
51 64.000000 64
52 64.000000 64
53 64.000000 64
54 64.000000 64
55 64.000000 64
56 64.000000 64
57 64.000000 64
58 64.000000 64
59 64.000000 64
60 64.000000 64
61 64.000000 64
62 64.000000 64
63 64.000000 64
$

Для новый код или новое обслуживание, вы должны использовать только версии _sync.

Дополнительные примеры использования параметра маски см. в этом блоге

4
ответ дан Robert Crovella 16 August 2018 в 02:06
поделиться
  • 1
    спасибо, это работает и для меня. Но каково значение маски «0xFFFFFFFF». в новой функции. – ztdep 2 June 2018 в 01:06
  • 2
    Вы прочитали документацию, на которую я указал? Например, этот абзац: «Новые * _sync shfl intrinsics принимают маску, указывающую потоки, участвующие в вызове. Бит, представляющий идентификатор полосы потока, должен быть установлен для каждого участвующего потока, чтобы убедиться, что они правильно конвергированы до того, как аппаратное обеспечение выполнено. Все не выходящие потоки, названные в маске, должны выполнять одно и то же внутреннее с той же маской, или результат не определен. & Quot; – Robert Crovella 2 June 2018 в 01:06
  • 3
    Я прочитал его. Но я не могу понять, как использовать маску и идентификатор полосы. – ztdep 2 June 2018 в 16:42
  • 4
    Если вы хотите эмулировать поведение старой инструкции non _sync, просто используйте 0xFFFFFFFF, как я показал. Если вы хотите больше узнать о новых инструкциях, это, вероятно, время для нового вопроса. Также есть блоги и презентации , которые обсуждают эту новую функцию в CUDA 9. Попробуйте прочитать блог, это информативно. – Robert Crovella 2 June 2018 в 18:00
  • 5
    В качестве альтернативы вы можете передать -1 маске для установки всех бит: temp+=__shfl_down_sync(-1, temp, offset,32); – AaronS 6 August 2018 в 16:37
Другие вопросы по тегам:

Похожие вопросы: